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
Opencl教程系列…
一，概述二，基本概念三，简单的demo
四，存储模型
……
请问这里的代码是在NvidIA的显卡上实现的，还是在AMD ATI显卡上实现的？