«
»

OpenCL系列讲座(三) Step By Step

1 Star2 Stars3 Stars4 Stars5 Stars
Loading ... Loading ...

  我们已经写好的Kernel函数一般保存在一个字符串内，例如在第二节中的kernel：输入两个一维数组，用另外一个一维数组输出两个一维数组求和的结果，如下：

char* chKernelSource = "__kernel sum 
(__global const float* a, __global const float* b,  
__global float* answer)
{
   int xid = get_global_id(0);
    answer[xid] = a[xid] + b[xid];
}";

如概述中所言，下面我们一步一步来写Host代码Main函数部分；
 
首先包含OpenCL的头文件

#include "cl/opencl.h";

 Step 1: 定义WorkItem和WorkGroup的大小：

const unsigned int cnGroupItemSize = 512; 
// 表示WorkGroup里的WorkItem以一维形式排列，一个WorkGroup里有512个WorkItem
 
const unsigned int cnGroupSize    = 3;
 // 表示WorkGroup以一维形式排列，一个Kernel里有3个WorkGroup
 
const unsigned int cnItemSize = cnGroupItemSize * cnGroupSize ; 
// 一个Kernel里运行的所有WorkItem的个数

扩展：

  • 如果WorkGroup和一个WorkGroup里的WorkItems以2维或3维形式组织的话，分别写成cnItemSize[i]和cnGroups[i]，i分别等于2和3；则cnItemSize等于 cnGroupItemSize[0] * cnGroupItemSize[1] *(cnGroupItemSize[2] )*cnGroups[0] * cnGroups[1] * (cnGroups[2]) 
  • cnItemSize是所有WorkItems的个数，和CUDA不同的是，CUDA用另外一个概念——GridDim来间接表示一个Kernel里面将运行多少个WorkItem（和CUDA里Thread的概念相同）。

Step 2: 创建显卡上下文(Device Context)

cl_context hContext; 
hContext = clCreateContextFromType(0, CL_DEVICE_TYPE_GPU, 
                                   0, 0, 0); // 指定由GPU显卡进行运算
size_t nContextDescriptorSize; 
 
clGetContextInfo(hContext, CL_CONTEXT_DEVICES,  
 
                 0, 0, &nContextDescriptorSize);  // 获得显卡ID字段的长度
 
cl_device_id * aDevices = malloc(nContextDescriptorSize); // 分配显卡ID字段的内存
 
clGetContextInfo(hContext, CL_CONTEXT_DEVICES, 
 
                 nContextDescriptorSize, aDevices, 0);  // 读取显卡ID字段
  • 简言之，显卡上下文就是对于运行Kernel代码的硬件的描述，存储了与硬件自身的参数，这里的hContext相当于一个handle，通过它可以得到安装的所有GPU硬件的信息
  • 由于OpenCL是介于硬件和软件开发人员之间的中间件，支持GPU/CPU/CELL，开发人员可以通过第二个参数，指定用何种计算显卡进行计算；

扩展：

  • 如果我们机器里面插了两块显卡，那么以上代码运行完后，将得到一个aDevices[]数组，存放了两块显卡的ID；
  • 这里通过clGetContextInfo函数，只得到了显卡ID这一个属性，如果我们想得到更多关于显卡的参数，可以照猫画虎，代码如下：

 

clGetContextInfo(hContext, CL_CONTEXT_PROPERTIES,  
 
                 0, 0, &nContextDescriptorSize);  // 获得显卡属性字段的长度
 
cl_context_properties * aDevicesProperty = malloc(nContextDescriptorSize); // 分配显卡属性字段的内存
 
clGetContextInfo(hContext, CL_CONTEXT_DEVICES, 
 
                 nContextDescriptorSize, aDevicesProperty , 0);  // 读取显卡属性字段
 
 

Step 3: 创建一个命令队列(Command Queue)

cl_command_queue hCmdQueue; 
 
hCmdQueue = clCreateCommandQueue(hContext, aDevices[0], 0, 0);  // 在查询到的0号显卡上创建命令队列
  • 命令队列用来排列（或存放）在Kernel的显卡上运行的一系列“命令”，如：依次运行不同的Kernel、Host和Device交换数据等。
  • Command Queue存在时有必要的：Host可以将在显卡上运行的命令依次“放到”这个队列里面，由队列来管理各命令的运行顺序（命令间可以是顺序的，也可以是异步的），而并不需要等待一个操作在显卡运行完毕，才返回；
  • Host可以进行不会对显卡运算产生冲突的其它操作，相当于在显卡上和在Host上运行不同的线程；
  • 在CUDA里，提供了一种简单的机制，默认在执行下一条命令前需要等待上一条命令执行完毕的；相比较，OpenCL引入的Command Queue机制更为灵活。

 

Step 4: 生成并加载BIN

cl_program hProgram; 
 
hProgram = clCreateProgramWithSource(hContext, 1,  
 
                                     &chKernelSource, 0, 0); // 从源代码创建一个适合于当前显卡上下文的cl_program对象
clBuildProgram(hProgram, 0, 0, 0, 0, 0);  // 编译此段代码
  • chKernelSource是一个字符串，里面包含了一个kernel代码的所有文字；
  • hProgram是一个句柄，指向加载的从源代码编译好的BIN代码（可执行代码）；
  • OpenCL本身不提供编译器，因此代码不会事先编译为能载入显卡的BIN代码，而是在运行时通过显卡厂商提供的编译器编译的；当然，OpenCL提供了加载并运行已生成的BIN代码的功能；
  • OpenCL作为一个中间层，感觉像是一个命令解释器，其实，在相同的显卡上，第二次加载BIN文件并不会重新编译，而是从缓存里面读取；

 

扩展

  • 可以将多段源码编译为BIN代码，只需要传入多个源码char*头指针即可
  • 用clCreateProgramWithBinary函数，可以跳过编译阶段，直接加载符合OpenCL规范的BIN代码

 

Step 5: 创建Kernel

cl_kernel hKernel; 
 
hKernel = clCreateKernel(hProgram, “vectorAdd”, 0); 
 
 
  • 创建Kernel的过程，可以看做从BIN代码提取Kernel对象的过程，需要指定Kernel函数名，且这个函数名在源代码中必须用__kernel关键字修饰；
  • clCreateKernel一次只提取一个Kernel对象，如果需要同时提取多个Kernel对象时，调用clCreateKernelsInProgram函数

 

Step 6: 在Host上分配内存，并读入数据

float * pA = new float[cnItemSize]; 
 
float * pB = new float[cnItemSize]; 
 
float * pC = new float[cnItemSize]; 
 
 

Step 7: 分配显卡内存，并传入数据

cl_mem hDeviceMemA, hDeviceMemB, hDeviceMemC; 
 
hDeviceMemA = clCreateBuffer(hContext,  
 
                          CL_MEM_READ_ONLY |   CL_MEM_COPY_HOST_PTR, 
 
                          cnItemSize * sizeof(cl_float),  
 
                          pA,  
 
                          0); 
 
hDeviceMemB = clCreateBuffer(hContext,  
 
                          CL_MEM_READ_ONLY |   CL_MEM_COPY_HOST_PTR, 
 
                          cnItemSize * sizeof(cl_float),  
 
                          pA,  
 
                          0); 
 
hDeviceMemC = clCreateBuffer(hContext,  
 
                          CL_MEM_WRITE_ONLY, 
 
                          cnItemSize * sizeof(cl_float),  
 
                          0, 0); 
  • OpenCL分配显卡内存，不返单纯返回显卡内存的指针，而是通过“Buffer Object”来管理；和CUDA比起来，优点是可以设置待分配内存的读写属性，还可以设置Host内存的指针，分配后直接从Host拷贝数据，不用再次调用写显卡内存函数。

 

扩展：

  • 读写“Buffer Object”用另外两个函数来实现，分别是：clEnqueueReadBuffer和clEnqueueWriteBuffer。
  • 在“Buffer Object”间拷贝数据用clEnqueueCopyBuffer函数。
  • clCreateBuffer函数中，第二个参数如果设置为CL_MEM_ALLOC_HOST_PTR，将分配Host内存给Kernel使用；这样可以以减少Host和Device交换数据的时间。Kernel调用完成后，执行clEnqueueMapBuffer函数，然后再从Host内存中读取数据。

 

Step 8: 为Kernel函数指定输入参数

clSetKernelArg(hKernel, 0, sizeof(cl_mem), (void *)&hDeviceMemA); 
 
clSetKernelArg(hKernel, 1, sizeof(cl_mem), (void *)&hDeviceMemB); 
 
clSetKernelArg(hKernel, 2, sizeof(cl_mem), (void *)&hDeviceMemC); 
  • 第二个参数表明输入参数和函数声明中的第几个参数相对应

 

Step 9: 执行Kernel函数

clEnqueueNDRangeKernel(hCmdQueue, hKernel, 1, 0,  
 
                       &cnGroupSize, &cnItemSize, 0); 

 

Step 10: 从显卡内存拷贝运行结果并释放内存

clEnqueueReadBuffer(hContext, hDeviceC, CL_TRUE, 0,  
 
                    cnItemSize * sizeof(cl_float),  
 
                    pC, 0, 0, 0); 
 
delete[] pA; 
 
delete[] pB; 
 
delete[] pC; 
 
clReleaseMemObj(hDeviceMemA); 
 
clReleaseMemObj(hDeviceMemB); 
 
clReleaseMemObj(hDeviceMemC); 

 

 

参考：

1. NVIDIA OpenCL JumpStart Guide 0.9

2. http://developer.apple.com/mac/library/documentation/Performance/Conceptual/
OpenCL_MacProgGuide/

3. http://www.macresearch.org/category/tutorials

4. http://www.khronos.org/message_boards/viewforum.php?f=28&sid=9f376ea1750bb6c84e8f361d5ff64cd0

标签： , | Print Print | 287 views

一条评论 发表在“OpenCL系列讲座(三) Step By Step”上

  1. GPGPU 说：
    2009年11月30日于11:27 上午

    Opencl教程系列…

    一，概述二，基本概念三，简单的demo
    四，存储模型
    ……

    回复

留下回复