OpenCL编程详细解析与实例
C语言与OpenCL的编程示例比较
参考链接：
https://www.zhihu.com/people/wujianming_110117/posts
先以图像旋转的实例，具体介绍OpenCL编程的步骤。 首先给出实现流程，然后给出实现图像旋转的C循环实现和OpenCL C kernel实现。
图像旋转原理
图像旋转是指把定义的图像绕某一点以逆时针或顺时针方向旋转一定的角度， 通常是指绕图像的中心以逆时针方向旋转。假设图像的左上角为(l, t), 右下角为(r, b)，则图像上任意点(x, y) 绕其中心(xcenter, ycenter)逆时针旋转θ角度后， 新的坐标位置(x',y')的计算公式为：
x′ = (x - xcenter) cosθ - (y － ycenter) sinθ + xcenter,
y′ = (x - xcenter) sinθ + (y － ycenter) cosθ + ycenter.
C代码：
void rotate(unsigned char* inbuf,unsigned char* outbuf,int w, int h,float sinTheta,float cosTheta){int i, j;int xc = w/2;int yc = h/2;for(i = 0; i < h; i++){for(j=0; j< w; j++){int xpos = (j-xc)*cosTheta - (i - yc) * sinTheta + xc;int ypos = (j-xc)*sinTheta + (i - yc) * cosTheta + yc;if(xpos>=0&&ypos>=0&&xpos<w&&ypos<h)outbuf[ypos*w + xpos] = inbuf[i*w+j];}}}OpenCL C kernel代码：
#pragma OPENCL EXTENSION cl_amd_printf : enable__kernel void image_rotate(__global uchar * src_data,__global uchar * dest_data, //Data in global memoryint W, int H, //Image Dimensionsfloat sinTheta, float cosTheta ) //Rotation Parameters{const int ix = get_global_id(0);const int iy = get_global_id(1);int xc = W/2;int yc = H/2;int xpos = ( ix-xc)*cosTheta - (iy-yc)*sinTheta+xc;int ypos = (ix-xc)*sinTheta + ( iy-yc)*cosTheta+yc;if ((xpos>=0) && (xpos< W) && (ypos>=0) && (ypos< H))dest_data[ypos*W+xpos]= src_data[iy*W+ix];}
正如上面代码中所给出的那样，在C代码中需要两重循环来计算横纵坐标上新的 坐标位置。其实，在图像旋转的算法中每个点的计算可以独立进行，与其它点的 坐标位置没有关系，所以并行处理较为方便。OpenCL C kernel代码中用了并行 处理。
上面的代码在Intel的OpenCL平台上进行了测试，处理器为双核处理器，图像大小 为4288*3216，如果用循环的方式运行时间稳定在0.256s左右，而如果用OpenCL C kernel并行的方式，运行时间稳定在0.132秒左右。GPU的测试在NVIDIA的GeForce G105M显卡 上进行，运行时间稳定在0.0810s左右。从循环的方式，双核CPU并行以及GPU并行计算 已经可以看出，OpenCL编程的确能大大提高执行效率。
OpenCL编程详细解析
OpenCL作为一门开源的异构并行计算语言，设计之初就是使用一种模型来模糊各种硬件差异。作为软件开发人员，关注的就是编程模型。OpenCL程序的流程大致如下：
- 加载 OpenCL 内核程序并创建一个 program 对象
- 为指定的 device 编译 program 中的 kernel
- 创建指定名字的 kernel 对象
- 为 kernel 创建内存对象
- 为 kernel 设置参数
- 在指定的 device 上创建 command queue
- 将要执行的 kernel 放入 command queue
- 将结果读回 host
- 资源回收
模块分析
使用 OpenCL API 编程与一般 C/C++ 引入第三方库编程没什么区别。所以，首先要做的自然是 include 相关的头文件。由于在 MacOS X 10.6下OpenCL的头文件命名与其他系统不同，通常使用一个#if defined进行区分，代码如下：
1. #if defined(__APPLE__) || defined(__MACOSX)2. #include <OpenCL/cl.hpp>3. #else4. #include <CL/cl.h>5. #endif接下来就进入真正的编码流程了。
Platform
查询并选择一个 platform
首先要取得系统中所有的 OpenCL platform。所谓的 platform 指的就是硬件厂商提供的 OpenCL 框架，不同的 CPU/GPU 开发商（比如 Intel、AMD、Nvdia）可以在一个系统上分别定义自己的 OpenCL 框架。所以需要查询系统中可用的 OpenCL 框架，即 platform。使用 API 函数 clGetPlatformIDs 获取可用 platform 的数量：
1. cl_int status = 0;2. cl_uint numPlatforms;3. cl_platform_id platform = NULL;4. status = clGetPlatformIDs( 0, NULL, &numPlatforms);5.6. if(status != CL_SUCCESS){7. printf("Error: Getting Platforms\n");8. return EXIT_FAILURE;9. }然后根据数量来分配内存，并得到所有可用的 platform，所使用的 API 还是clGetPlatformIDs。在 OpenCL 中，类似这样的函数调用很常见：第一次调用以取得数目，便于分配足够的内存；然后调用第二次以获取真正的信息。
1. if (numPlatforms > 0) {2. cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));3. status = clGetPlatformIDs(numPlatforms, platforms, NULL);4. if (status != CL_SUCCESS) {5. printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");6. return -1;7. }现在，所有的 platform 都存在了变量 platforms 中，接下来需要做的就是取得所需的 platform。本人的PC上配置的是 Intel 处理器和 AMD 显卡，专业点的说法叫 Intel 的 CPU 和 NVIDIA的 GPU :)。所以这儿有两套 platform，为了体验下 GPU 的快感，所以使用 AMD 的 platform。通过使用 clGetPlatformInfo 来获得 platform 的信息。通过这个 API 可以知晓 platform 的厂商信息，以便选出需要的 platform。代码如下：
1. for (unsigned int i = 0; i < numPlatforms; ++i) {2. char pbuff[100];3. status = clGetPlatformInfo(4. platforms[i],5. CL_PLATFORM_VENDOR,6. sizeof(pbuff),7. pbuff,8. NULL);9. platform = platforms[i];10. if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {11. break;12. }13. }不同的厂商信息可以参考 OpenCL Specifications，这儿只是简单的筛选出 AMD 。
在 platform 上建立 context
第一步是通过 platform 得到相应的 context properties
1. // 如果能找到相应平台，就使用，否则返回NULL2. cl_context_properties cps[3] = {3. CL_CONTEXT_PLATFORM,4. (cl_context_properties)platform,5. 06. };7.8. cl_context_properties *cprops = (NULL == platform) ? NULL : cps;第二步是通过 clCreateContextFromType 函数创建 context。
1. // 生成 context2. cl_context context = clCreateContextFromType(3. cprops,4. CL_DEVICE_TYPE_GPU,5. NULL,6. NULL,7. &status);8. if (status != CL_SUCCESS) {9. printf("Error: Creating Context.(clCreateContexFromType)\n");10. return EXIT_FAILURE;11. }函数的第二个参数可以设定 context 关联的设备类型。本例使用的是 GPU 作为OpenCL计算设备。目前可以使用的类别包括：
1. - CL_DEVICE_TYPE_CPU2. - CL_DEVICE_TYPE_GPU3. - CL_DEVICE_TYPE_ACCELERATOR4. - CL_DEVICE_TYPE_DEFAULT5. - CL_DEVICE_TYPE_ALL在 context 上查询 device
context 创建好之后，要做的就是查询可用的 device。
1. status = clGetContextInfo(context,2. CL_CONTEXT_DEVICES,3. 0,4. NULL,5. &deviceListSize);6. if (status != CL_SUCCESS) {7. printf("Error: Getting Context Info device list size, clGetContextInfo)\n");8. return EXIT_FAILURE;9. }10. cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);11. if (devices == 0) {12. printf("Error: No devices found.\n");13. return EXIT_FAILURE;14. }15.16. status = clGetContextInfo(context,17. CL_CONTEXT_DEVICES,18. deviceListSize,19. devices,20. NULL);21. if (status != CL_SUCCESS) {22. printf("Error: Getting Context Info (device list, clGetContextInfo)\n");23. return EXIT_FAILURE;24. }与获取 platform 类似，调用两次 clGetContextInfo 来完成 查询。第一次调用获取关联 context 的 device 个数，并根据个数申请内存；第二次调用获取所有 device 实例。如果想了解每个 device 的具体信息，可以调用 clGetDeviceInfo 函数来获取，返回的信息有设备类型、生产商以及设备对某些扩展功能的支持与否等等。详细使用情况请参阅 OpenCL Specifications。
到此，platform 相关的程序已经准备就绪了，下面到此的完整代码：
1. /* OpenCL_01.cpp2. * (c) by keyring <keyrings@163.com>3. * 2013.10.264. */5.6. #if defined(__APPLE__) || defined(__MACOSX)7. #include <OpenCL/cl.hpp>8. #else9. #include <CL/cl.h>10. #endif11.12. #include <iostream>13.14. int main(int argc, char const *argv[])15. {16. printf("hello OpenCL\n");17. cl_int status = 0;18. size_t deviceListSize;19.20. // 得到并选择可用平台21. cl_uint numPlatforms;22. cl_platform_id platform = NULL;23. status = clGetPlatformIDs(0, NULL, &numPlatforms);24.25. if (status != CL_SUCCESS) {26. printf("ERROR: Getting Platforms.(clGetPlatformIDs)\n");27. return EXIT_FAILURE;28. }29.30. if (numPlatforms > 0) {31. cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));32. status = clGetPlatformIDs(numPlatforms, platforms, NULL);33. if (status != CL_SUCCESS) {34. printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");35. return -1;36. }37.38. // 遍历所有 platform，选择想用的39. for (unsigned int i = 0; i < numPlatforms; ++i) {40. char pbuff[100];41. status = clGetPlatformInfo(42. platforms[i],43. CL_PLATFORM_VENDOR,44. sizeof(pbuff),45. pbuff,46. NULL);47. platform = platforms[i];48. if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {49. break;50. }51. }52.53. delete platforms;54. }55.56. // 如果能找到相应平台，就使用，否则返回NULL57. cl_context_properties cps[3] = {58. CL_CONTEXT_PLATFORM,59. (cl_context_properties)platform,60. 061. };62.63. cl_context_properties *cprops = (NULL == platform) ? NULL : cps;64.65.66. // 生成 context67. cl_context context = clCreateContextFromType(68. cprops,69. CL_DEVICE_TYPE_GPU,70. NULL,71. NULL,72. &status);73. if (status != CL_SUCCESS) {74. printf("Error: Creating Context.(clCreateContexFromType)\n");75. return EXIT_FAILURE;76. }77.78. // 寻找OpenCL设备79.80. // 首先得到设备列表的长度81. status = clGetContextInfo(context,82. CL_CONTEXT_DEVICES,83. 0,84. NULL,85. &deviceListSize);86. if (status != CL_SUCCESS) {87. printf("Error: Getting Context Info device list size, clGetContextInfo)\n");88. return EXIT_FAILURE;89. }90. cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);91. if (devices == 0) {92. printf("Error: No devices found.\n");93. return EXIT_FAILURE;94. }95.96. // 然后得到设备列表97. status = clGetContextInfo(context,98. CL_CONTEXT_DEVICES,99. deviceListSize,100. devices,101. NULL);102. if (status != CL_SUCCESS) {103. printf("Error: Getting Context Info (device list, clGetContextInfo)\n");104. return EXIT_FAILURE;105. }Running time
前面写了这么多，其实还没真正进入具体的程序逻辑中，顶多算配好了 OpenCL 运行环境。真正的逻辑代码，即程序的任务就是运行时模块。本例的任务是在一个 4×4的二维空间上，按一定的规则给每个元素赋值，具体代码如下：
1. #define KERNEL(...)#__VA_ARGS__2.3. const char *kernelSourceCode = KERNEL(4. __kernel void hellocl(__global uint *buffer)5. {6. size_t gidx = get_global_id(0);7. size_t gidy = get_global_id(1);8. size_t lidx = get_local_id(0);9. buffer[gidx + 4 * gidy] = (1 << gidx) | (0x10 << gidy);10.11. }12. );这一段就是真正的逻辑，也就是代码要干的事。使用的是 OpenCL 自定的一门类C语言，具体的语法什么的现在先不纠结。这段代码是直接嵌入 cpp 文件的静态字符串。也可以将 kernel 程序单独写成一个文件。
加载 OpenCL 内核程序并创建一个 program 对象
接下来要做的就是读入 OpenCL kernel 程序并创建一个 program 对象。
1. size_t sourceSize[] = {strlen(kernelSourceCode)};2. cl_program program = clCreateProgramWithSource(context,3. 1,4. &kernelSourceCode,5. sourceSize,6. &status);7. if (status != CL_SUCCESS) {8. printf("Error: Loading Binary into cl_program (clCreateProgramWithBinary)\n");9. return EXIT_FAILURE;10. }本例中的 kernel 程序是作为静态字符串读入的（单独的文本文件也一样），所以使用的是 clCreateProgramWithSource，如果不想让 kernel 程序让其他人看见，可以先生成二进制文件，再通过 clCreateProgramWithBinary 函数动态读入二进制文件，做一定的保密。详细请参阅 OpenCL Specifications。
为指定的 device 编译 program 中的 kernel
kernel 程序读入完毕，要做的自然是使用 clBuildProgram 编译 kernel：
1. status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);2. if (status != CL_SUCCESS) {3. printf("Error: Building Program (clBuildingProgram)\n");4. return EXIT_FAILURE;5. }最终，kernel 将被相应 device 上的 OpenCL 编译器编译成可执行的机器码。
创建指定名字的 kernel 对象
成功编译后，可以通过 clCreateKernel 来创建一个 kernel 对象。
1. cl_kernel kernel = clCreateKernel(program, "hellocl", &status);2. if (status != CL_SUCCESS) {3. printf("Error: Creating Kernel from program.(clCreateKernel)\n");4. return EXIT_FAILURE;5. }引号中的 hellocl 就是 kernel 对象所关联的 kernel 函数的函数名。要注意的是，每个 kernel 对象必须关联且只能关联一个包含于相应 program 对象内的 kernel 程序。实际上，用户可以在 cl 源代码中写任意多个 kernel 程序，但在执行某个 kernel 程序之前必须先建立单独的 kernel 对象，即多次调用 clCreateKernel 函数。
为 kernel 创建内存对象
OpenCL 内存对象是指在 host 中创建，用于 kernel 程序的内存类型。按维度可以分为两类，一类是 buffer，一类是 image。buffer 是一维的，image 可以是二维、三维的 texture、frame-buffer 或 image。本例仅仅使用 buffer，可以通过clCreateBuffer 函数来创建。
1. cl_mem outputBuffer = clCreateBuffer(2. context,3. CL_MEM_ALLOC_HOST_PTR,4. 4 * 4 * 4,5. NULL,6. &status);7. if (status != CL_SUCCESS) {8. printf("Error: Create Buffer, outputBuffer. (clCreateBuffer)\n");9. return EXIT_FAILURE;10. }为 kernel 设置参数
使用 clSetKernelArg 函数为 kernel 设置参数。传递的参数既可以是常数，变量，也可以是内存对象。本例传递的就是内存对象。
1. status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&outputBuffer);2. if (status != CL_SUCCESS) {3. printf("Error: Setting kernel argument. (clSetKernelArg)\n");4. return EXIT_FAILURE;5. }该函数每次只能设置一个参数，如有多个参数，需多次调用。而且 kernel 程序中所有的参数都必须被设置，否则在启动 kernel 程序是会报错。指定位置的参数的类型最好和对应 kernel 函数内参数类型一致，以免产生各种未知的错误。在设置好指定参数后，每次运行该 kernel 程序都会使用设置值，直到用户使用次 API 重新设置参数。
在指定的 device 上创建 command queue
command queue 用于光里将要执行的各种命令。可以通过 clCreateCommandQueue函数创建。其中的 device 必须为 context 的关联设备，所有该 command queue 中的命令都会在这个指定的 device 上运行。
1. cl_command_queue commandQueue = clCreateCommandQueue(context,2. devices[0],3. 0,4. &status);5. if (status != CL_SUCCESS) {6. printf("Error: Create Command Queue. (clCreateCommandQueue)\n");7. return EXIT_FAILURE;8. }将要执行的 kernel 放入 command queue
创建好 command queue 后，用户可以创建相应的命令并放入 command queue 中执行。OpenCL 提供了三种方案来创建 kernel 执行命令。最常用的即为本例所示的运行在指定工作空间上的 kernel 程序，使用了 clEnqueueNDRangeKernel 函数。
1. size_t globalThreads[] = {4, 4};2. size_t localThreads[] = {2, 2};3. status = clEnqueueNDRangeKernel(commandQueue, kernel,4. 2, NULL, globalThreads,5. localThreads, 0,6. NULL, NULL);7. if (status != CL_SUCCESS) {8. printf("Error: Enqueueing kernel\n");9. return EXIT_FAILURE;10. }clEnqueueNDRangeKernel 函数每次只能将一个 kernel 对象放入 command queue 中，用户可以多次调用该 API 将多个 kernel 对象放置到一个 command queue 中，command queue 中的不同 kernel 对象的工作区域完全不相关。其余两个 APIclEnqueueTask 和 clEnqueueNativeKernel 的用法就不多讲了，详情请参阅OpenCL Specificarions。
最后可以用 clFinish 函数来确认一个 command queue 中所有的命令都执行完毕。函数会在 command queue 中所有 kernel 执行完毕后返回。
1. // 确认 command queue 中所有命令都执行完毕2. status = clFinish(commandQueue);3. if (status != CL_SUCCESS) {4. printf("Error: Finish command queue\n");5. return EXIT_FAILURE;6. }将结果读回 host
计算完毕，将结果读回 host 端。使用 clEnqueueReadBuffer 函数将 OpenCL buffer 对象中的内容读取到 host 可以访问的内存空间。
1. // 将内存对象中的结果读回Host2. status = clEnqueueReadBuffer(commandQueue,3. outputBuffer, CL_TRUE, 0,4. 4 * 4 * 4, outbuffer, 0, NULL, NULL);5. if (status != CL_SUCCESS) {6. printf("Error: Read buffer queue\n");7. return EXIT_FAILURE;8. }当然，为了看下程序的运行效果，咱们当然得看看运行结果啦。打印一下吧：
1. // Host端打印结果2. printf("out:\n");3. for (int i = 0; i < 16; ++i) {4. printf("%x ", outbuffer[i]);5. if ((i + 1) % 4 == 0)6. printf("\n");7. }资源回收
程序的最后是对所有创建的对象进行释放回收，与C/C++的内存回收同理。
1. // 资源回收2. status = clReleaseKernel(kernel);3. status = clReleaseProgram(program);4. status = clReleaseMemObject(outputBuffer);5. status = clReleaseCommandQueue(commandQueue);6. status = clReleaseContext(context);7.8. free(devices);9. delete outbuffer;总结
这次使用一个小例子来详细说明了 OpenCL 编程的一般步骤。其实这些步骤一般都是固定的。真正需要注意的是 OpenCL Kernel 程序的编写。当然，合理高效的利用 API 也是一门技术活。
完整程序
1. #include <iostream>2. #include <stdlib.h>3. #include <string.h>4. #include <stdio.h>5.6.7. #if defined(__APPLE__) || defined(__MACOSX)8. #include <OpenCL/cl.hpp>9. #else10. #include <CL/cl.h>11. #endif12.13. using namespace std;14.15. #define KERNEL(...)#__VA_ARGS__16.17. const char *kernelSourceCode = KERNEL(18. __kernel void hellocl(__global uint *buffer)19. {20. size_t gidx = get_global_id(0);21. size_t gidy = get_global_id(1);22. size_t lidx = get_local_id(0);23. buffer[gidx + 4 * gidy] = (1 << gidx) | (0x10 << gidy);24.25. }26. );27.28. int main(int argc, char const *argv[])29. {30. printf("hello OpenCL\n");31. cl_int status = 0;32. size_t deviceListSize;33.34. // 当前服务器上配置的仅有NVIDIA Tesla C2050 的GPU35. cl_platform_id platform = NULL;36. status = clGetPlatformIDs(1, &platform, NULL);37.38. if (status != CL_SUCCESS) {39. printf("ERROR: Getting Platforms.(clGetPlatformIDs)\n");40. return EXIT_FAILURE;41. }42.43. // 如果能找到相应平台，就使用，否则返回NULL44. cl_context_properties cps[3] = {45. CL_CONTEXT_PLATFORM,46. (cl_context_properties)platform,47. 048. };49.50. cl_context_properties *cprops = (NULL == platform) ? NULL : cps;51.52.53. // 生成 context54. cl_context context = clCreateContextFromType(55. cprops,56. CL_DEVICE_TYPE_GPU,57. NULL,58. NULL,59. &status);60. if (status != CL_SUCCESS) {61. printf("Error: Creating Context.(clCreateContexFromType)\n");62. return EXIT_FAILURE;63. }64.65. // 寻找OpenCL设备66.67. // 首先得到设备列表的长度68. status = clGetContextInfo(context,69. CL_CONTEXT_DEVICES,70. 0,71. NULL,72. &deviceListSize);73. if (status != CL_SUCCESS) {74. printf("Error: Getting Context Info device list size, clGetContextInfo)\n");75. return EXIT_FAILURE;76. }77. cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);78. if (devices == 0) {79. printf("Error: No devices found.\n");80. return EXIT_FAILURE;81. }82.83. // 现在得到设备列表84. status = clGetContextInfo(context,85. CL_CONTEXT_DEVICES,86. deviceListSize,87. devices,88. NULL);89. if (status != CL_SUCCESS) {90. printf("Error: Getting Context Info (device list, clGetContextInfo)\n");91. return EXIT_FAILURE;92. }93.94.95. // 装载内核程序，编译CL program ,生成CL内核实例96.97. size_t sourceSize[] = {strlen(kernelSourceCode)};98. cl_program program = clCreateProgramWithSource(context,99. 1,100. &kernelSourceCode,101. sourceSize,102. &status);103. if (status != CL_SUCCESS) {104. printf("Error: Loading Binary into cl_program (clCreateProgramWithBinary)\n");105. return EXIT_FAILURE;106. }107.108. // 为指定的设备编译CL program.109. status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);110. if (status != CL_SUCCESS) {111. printf("Error: Building Program (clBuildingProgram)\n");112. return EXIT_FAILURE;113. }114.115. // 得到指定名字的内核实例的句柄116. cl_kernel kernel = clCreateKernel(program, "hellocl", &status);117. if (status != CL_SUCCESS) {118. printf("Error: Creating Kernel from program.(clCreateKernel)\n");119. return EXIT_FAILURE;120. }121.122. // 创建 OpenCL buffer 对象123. unsigned int *outbuffer = new unsigned int [4 * 4];124. memset(outbuffer, 0, 4 * 4 * 4);125. cl_mem outputBuffer = clCreateBuffer(126. context,127. CL_MEM_ALLOC_HOST_PTR,128. 4 * 4 * 4,129. NULL,130. &status);131.132. if (status != CL_SUCCESS) {133. printf("Error: Create Buffer, outputBuffer. (clCreateBuffer)\n");134. return EXIT_FAILURE;135. }136.137.138. // 为内核程序设置参数139. status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&outputBuffer);140. if (status != CL_SUCCESS) {141. printf("Error: Setting kernel argument. (clSetKernelArg)\n");142. return EXIT_FAILURE;143. }144.145. // 创建一个OpenCL command queue146. cl_command_queue commandQueue = clCreateCommandQueue(context,147. devices[0],148. 0,149. &status);150. if (status != CL_SUCCESS) {151. printf("Error: Create Command Queue. (clCreateCommandQueue)\n");152. return EXIT_FAILURE;153. }154.155.156. // 将一个kernel 放入 command queue157. size_t globalThreads[] = {4, 4};158. size_t localThreads[] = {2, 2};159. status = clEnqueueNDRangeKernel(commandQueue, kernel,160. 2, NULL, globalThreads,161. localThreads, 0,162. NULL, NULL);163. if (status != CL_SUCCESS) {164. printf("Error: Enqueueing kernel\n");165. return EXIT_FAILURE;166. }167.168. // 确认 command queue 中所有命令都执行完毕169. status = clFinish(commandQueue);170. if (status != CL_SUCCESS) {171. printf("Error: Finish command queue\n");172. return EXIT_FAILURE;173. }174.175. // 将内存对象中的结果读回Host176. status = clEnqueueReadBuffer(commandQueue,177. outputBuffer, CL_TRUE, 0,178. 4 * 4 * 4, outbuffer, 0, NULL, NULL);179. if (status != CL_SUCCESS) {180. printf("Error: Read buffer queue\n");181. return EXIT_FAILURE;182. }183.184. // Host端打印结果185. printf("out:\n");186. for (int i = 0; i < 16; ++i) {187. printf("%x ", outbuffer[i]);188. if ((i + 1) % 4 == 0)189. printf("\n");190. }191.192. // 资源回收193. status = clReleaseKernel(kernel);194. status = clReleaseProgram(program);195. status = clReleaseMemObject(outputBuffer);196. status = clReleaseCommandQueue(commandQueue);197. status = clReleaseContext(context);198.199. free(devices);200. delete outbuffer;201. return 0;202. }运行结果
参考链接：
https://www.zhihu.com/people/wujianming_110117/posts
人工智能芯片与自动驾驶
