3.6 OpenCL运行时(例子)

OpenCL的四种模型在之前的章节中已经全部讨论了，OpenCL通过运行时API让应用开发者了解这些模型。平台模型用来使能一个主机，以及一个或多个设备，让其参与到OpenCL应用的执行中。应用开发者使用编程模型来让OpenCL内核实现其核心计算部分。内核执行时如何获取需要的数据，则有内存模型定义。开发者通过执行模型提交相应的命令到设备端(执行内存搬运或执行内核任务)。本节会将这些内容融合到一个完整的OpenCL应用中。

创建并执行一个简单的OpenCL应用大致需要以下几步：

  1. 查询平台和设备信息

  2. 创建一个上下文

  3. 为每个设备创建一个命令队列

  4. 创建一个内存对象(数组)用于存储数据

  5. 拷贝输入数据到设备端

  6. 使用OpenCL C代码创建并编译出一个程序

  7. 从编译好的OpenCL程序中提取内核

  8. 执行内核

  9. 拷贝输出数据到主机端

  10. 释放资源

下面的代码将具体实现以上总结的每一步。OpenCL应用中，大部分的通用代码对OpenCL的执行进行设置，这样就允许跨硬件平台(不同的供应商)架构来执行OpenCL内核。因此，其中的大部分代码可以直接在其他的应用中直接使用，并且可以抽象成用户定义函数。之后的章节，将展示OpenCL C++ API，其冗余要小于C API。

现在我们来讨论逐个步骤。本节之后，将会提供完整的程序代码。

1. 查询平台和设备

OpenCL内核需要执行在设备端，那么就需要至少一个平台和一个设备可以查询。

  1. cl_int status; // 用于错误检查
  3. // 检索平台的数量
  4. cl_uint numPlatforms = 0;
  5. status = clGetPlatformIDs(0, NULL, &numPlatforms);
  7. // 为每个平台对象分配足够的空间
  8. cl_platform_id *platforms = NULL;
  9. platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));
  11. // 将具体的平台对象填充其中
  12. status = clGetPlatformIDs(numPlatforms, platforms, NULL);
  14. // 检索设备的数量
  15. cl_uint numDevices = 0;
  16. status = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, 0, NULL, &numDevices);
  18. // 为每个设备对象分配足够的空间
  19. cl_device_id *devices;
  20. devices = (cl_device_id *)malloc(numDevices * sizeof(cl_device_id));
  22. // 将具体的设备对象填充其中
  23. status = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, numDevices, devices, NULL);

之后的完整代码中，我们将默认使用首先找到的平台和设备。这样的源码看起来更加简单明了。

2. 创建一个上下文

找到平台和设备之后，就可以在主机端对上下文进行配置。

  1. // 创建的上下文包含所有找到的设备
  2. cl_context context = clCreateContext(NULL, numDevices, devices, NULL, NULL, &status);

3. 为每个设备创建一个命令队列

创建完上下文，就要为每个设备创建一个命令队列(每个命令队列只关联其对应的设备)。主机端需要设备端执行的命令将提交到命令队列中，由命令队列管理执行。

  1. // 为第一个发现的设备创建命令队列
  2. cl_command_queue cmdQueue = clCreateCommandQUeueWithProperties(context, devices[0], 0, &status);

4. 创建设备数组用于存储数据

创建一个数组需要提供其长度，以及与该数组相关的上下文；该数组能对该上下文所有设备可见。通常，调用者需要提供一些标识，来表明数据是可只读、只写或读写。如果第四个参数传NULL，OpenCL将不会在这步对数组进行初始化。

  1. // 向量加法的三个向量，2个输入数组和1个输出数组
  2. cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status);
  3. cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status);
  4. cl_mem bufC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, datasize, NULL, &status);

5. 拷贝输入数据到设备端

下一步就是将主机端指针指向的数组拷贝到设备端。该API需要一个命令队列对象作为参数，所以数据通常都是直接拷贝到设备端。将第三个参数设置为CL_TRUE，我们将等待该API将数据全部拷贝到设备端之后才返回。

  1. // 将输入数据填充到数组中
  2. status = clEnqueueWriteBuffer(cmdQueue, bufA, CL_TRUE, 0, datasize, A, 0, NULL, NULL);
  3. status = clEnqueueWriteBuffer(cmdQueue, bufB, CL_TRUE, 0, datasize, B, 0, NULL, NULL);

6. 使用OpenCL C代码创建并编译出一个程序

代码列表3.3中的向量相加内核存储在一个字符数组中，programSource，当可以用来创建程序对象(之后需要编译)。当我们编译一个程序时，我们需要提供目标设备的信息。

  1. // 使用源码创建程序
  2. cl_program program = clCreateProgramWithSource(context, 1, (const char **)&programSource, NULL, &status);
  4. // 为设备构建(编译)程序
  5. status = clBuildProgram(program, numDevices, devices, NULL, NULL, NULL);

7. 从编译好的OpenCL程序中提取内核

内核通过提供内核函数名，在程序上进行创建。

  1. // 创建向量相加内核
  2. cl_kernel kernel = clCreateKernel(program, "vecadd", &status);

8. 执行内核

内核创建完毕，数据都已经传输到设备端，数组需要设置到内核的参数上。之后，一条需要执行内核的命令就进入了命令队列。内核的执行方式需要指定的NDRange进行配置。

  1. // 设置内核参数
  2. status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA);
  3. status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufB);
  4. status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufC);
  6. // 定义工作项的空间维度和空间大小
  7. // 虽然工作组的设置不是必须的，不过可以设置一下
  8. size_t indexSpaceSize[1],workGroupSize[1];
  10. indexSpaceSize[0] = datasize / sizeof(int);
  11. workGroupSize[0] = 256;
  13. // 通过执行API执行内核
  14. status = clEnqueueNDRangeKernel(cmdQueue, kernel, 1, NULL, indexSpaceSize, workGroupSize, 0, NULL, NULL);

9. 拷贝输出数据到主机端

  1. // 将输出数组拷贝到主机端内存中
  2. status = clEnqueueReadBuffer(cmdQueue, bufC, CL_TRUE, 0, datasize, C, 0, NULL, NULL);

10. 释放资源

内核执行完成后，并且输出已经传出到主机端，OpenCL分配的资源需要进行释放。这点和C/C++程序中的内存操作、文件处理以及其他资源的处理，都需要开发者显式释放。OpenCL为不同的对象提供了不同的释放API。OpenCL上下文需要最后释放，因为数组和命令队列都绑定在上下文上。这点与C++删除对象有些相似，成员数组需要在成员释放前释放。

  1. clReleaseKernel(kernel);
  2. clReleaseProgram(program);
  3. clReleaseCommandQueue(cmdQueue);
  4. clReleaseMemObject(bufA);
  5. clReleaseMemObject(bufB);
  6. clReleaseMemObject(bufC);
  7. clReleaseContext(context);

3.6.1 向量相加的完整代码

下面将完整的展示向量相加这个例子。其具有上一节的所有步骤，不过这个例子中使用了第一个平台对象和设备对象。

  1. // This program implements a vector addition using OpenCL
  3. // System includes
  4. #include <stdio.h>
  5. #include <stdlin.h>
  6. // OpenCL includes
  7. #include <CL/cl.h>
  9. // OpenCL kernel to perform an element-wise addition
  10. const char *programSouce = 
  11. "__kernel                                         \n"
  12. "void vecadd(__global int *A,                     \n"
  13. "            __global int *B,                     \n"
  14. "            __global int *C)                     \n"
  15. "{                                                \n"
  16. "  // Get the work-item's unique ID               \n"
  17. "  int idx = get_global_id(0);                    \n"
  18. "                                                 \n"
  19. "  // Add the corresponding locations of          \n"
  20. "  // 'A' and 'B', and store the reasult in 'C'   \n"
  21. "  C[idx] = A[idx] + B[idx];                      \n"
  22. "}                                                \n"
  23. ;
  25. int main(){
  26.   // This code executes on the OpenCL host
  28.   // Elements in each array
  29.   const int elements = 2048;
  31.   // Compute the size of the data
  32.   size_t datasize = sizeof(int) * elements;
  34.   // Allocate space for input/output host data
  35.   int *A = (int *)malloc(datasize); // Input array
  36.   int *B = (int *)malloc(datasize); // Input array
  37.   int *C = (int *)malloc(datasize); // Output array
  39.   // Initialize the input data
  40.   int i;
  41.   for (i = 0; i < elements; i++){
  42.     A[i] = i;
  43.     B[i] = i;
  44.   }
  46.   // Use this to check the output of each API call
  47.   cl_int status;
  49.   // Get the first platforms
  50.   cl_platform_id platform;
  51.   status = clGetPlatformIDs(1, &perform, NULL);
  53.   // Get the first devices
  54.   cl_device_id device;
  55.   status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, NULL);
  57.   // Create a context and associate it with the device
  58.   cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, &status);
  60.   // Create a command-queue and associate it with device
  61.   cl_command_queue cmdQueue = clCreateCommandQUeueWithProperties(context, device, 0, &status);
  63.   // Allocate two input buffers and one output buffer for the three vectors in the vector addition
  64.   cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status);
  65.   cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status);
  66.   cl_mem bufC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, datasize, NULL, &status);
  68.   // Write data from the input arrays to the buffers
  69.   status = clEnqueueWriteBuffer(cmdQueue, bufA, CL_FALSE, 0, datasize, A, 0, NULL, NULL);
  70.   status = clEnqueueWriteBuffer(cmdQueue, bufB, CL_FALSE, 0, datasize, B, 0, NULL, NULL);
  72.   // Create a program with source code
  73.   cl_program program = clCreateProgramWithSource(context, 1, (const char**)&programSource, NULL, &status);
  75.   // Build(compile) the program for the device
  76.   status = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
  78.   // Create the vector addition kernel
  79.   cl_kernel kernel = clCreateKernel(program, "vecadd", &status);
  81.   // Set the kernel arguments
  82.   status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA);
  83.   status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufB);
  84.   status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufC);
  86.   // Define an incde space of work-items for execution
  87.   // A work-group size is not required, but can be used.
  88.   size_t indexSpaceSize[1], workGroupSize[1];
  90.   // There are 'elements' work-items
  91.   indexSpaceSize[0] = elements;
  92.   workGroupSize[0] = 256;
  94.   // Execute the kernel
  95.   status = clEnqueueNDRangeKernel(cmdQueue, kernel, 1, NULL, indexSpaceSize, workGroupSize, 0, NULL, NULL);
  97.   // Read the device output buffer to the host output array
  98.   status = clEnqueueReadBuffer(cmdQueue, bufC, CL_TRUE, 0, datasize, C, 0, NULL, NULL);
  100.   // Free OpenCL resouces
  101.   clReleaseKernel(kernel);
  102.   clReleaseProgram(program);
  103.   clReleaseCommandQueue(cmdQueue);
  104.   clReleaseMemObject(bufA);
  105.   clReleaseMemObject(bufB);
  106.   clReleaseMemObject(bufC);
  107.   clReleaseContext(context);
  109.   // free host resouces
  110.   free(A);
  111.   free(B);
  112.   free(C);
  114.   return 0;
  115. }

代码清单3.4 使用C API实现的OpenCL向量相加