从 C 到 OpenCL

如何从 C/C++ 的 for 循环迁移到 OpenCL 内核。

算法

本教程中我们将探讨一个非常简单的算法。逐个元素相加两个数字阵列，并将结果存储在第三个阵列中：

C_n = A_n + B_n

虽然本算法非常简单，但它可从 OpenCL 实施中获益。阵列中的各个元素可以单独计算，因为阵列中元素之间没有依赖关系。这表示可以轻松地并行计算各个元素。对于这样的工作负载，OpenCL 是理想之选。

C/C++ 实施

假设我们有以下三个阵列：

 

1 2 3 4 5 6 7

/* Number of elements in the arrays of input and output data. */ int arraySize = 1000000;   /* Arrays to hold the input and output data. */ int* inputA = new int[arraySize]; int* inputB = new int[arraySize]; int* output = new int[arraySize];

 

那么，C/C++ 中的实施是琐碎的：

 

1 2 3 4

for (int i = 0; i < arraySize; i++) { output[i] = inputA[i] + inputB[i]; }

不计入任何编译器优化的话，这一代码将在 CPU 上顺序执行。这表示阵列中的各个元素将以串行方式计算。各个计算之间创建了一个人为的依赖关系。

此代码的运行时间与阵列的大小成比例。

OpenCL 实施

基本信息

1. 将可并行处理的代码迁移到 OpenCL 内核中 提取代码的并行部分（我们的示例中为 for 循环），将它移到 OpenCL 内核中。对于我们的代码，不进行任何优化

1.  

   1 2 3 4 5 6 7 8 9 10 11 12 13

   __kernel void hello_world_opencl(__global int* restrict inputA,                                  __global int* restrict inputB,                                  __global int* restrict output) {     /*      * Set i to be the ID of the kernel instance.      * If the global work size (set by clEnqueueNDRangeKernel) is n,      * then n kernels will be run and i will be in the range [0, n - 1].      */     int i = get_global_id(0);     /* Use i as an index into the three arrays. */     output[i] = inputA[i] + inputB[i]; }

    

2. 运行内核的多个实例 内核中没有循环；因此，如果要使该代码运算阵列的所有元素，我们必须运行同一内核的多个实例：

    

   1 2 3 4 5 6 7 8 9 10 11 12

   /*      * Each instance of our OpenCL kernel operates on a single element of each array so the number of      * instances needed is the number of elements in the array.      */     size_t globalWorksize[1] = {arraySize};     /* Enqueue the kernel */     if (!checkSuccess(clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, globalWorksize, NULL, 0, NULL, &event)))     {         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);         cerr << "Failed enqueuing the kernel. " << __FILE__ << ":"<< __LINE__ << endl;         return 1;     }

    

    

   这会将 arraySize 个内核实例提交到 OpenCL 设备。

   每个实例被分配一个唯一 ID，我们使用该 ID 来选择各个实例所运算的阵列元素（见上方的内核）。

   由于我们没有在内核之间指定任何依赖关系，OpenCL 设备将自由地并行运行内核的这些实例。现在并行处理的唯一限制在设备能力上。

   此代码的运行时与阵列大小除以可并行运行的内核实例数所得的结果成比例。

   OpenCL 设置

   要运行上述代码，首先需要对 OpenCL 进行一些设置。详情可参阅 SDK 包中的 hello_world_opencl.cpp。

   内存

   由于现在运算是在 GPU 而不是 CPU 上执行的，我们需要了解我们使用的任何数据的位置。务必要知道数据是 GPU 还是 CPU 内存空间。

   在桌面系统中，GPU 和 CPU 拥有其自己的内存，通过相对较慢的总线分隔。这意味着在 CPU 和 GPU 之间共享内容可能是代价昂贵的运算。

   在配有 Mali-T600 系列 GPU 的大多数嵌入式系统中，CPU 和 GPU 共享通用内存。因此，在 CPU 和 GPU 之间共享内存的代价可能相对较低。

   由于这些系统差异，OpenCL 支持通过多种方式在设备之间分配和共享内存。

   以下是一种在设备间共享内存的方式，它旨在免除从一个设备到另一个设备（共享内存系统中）的内存复制：

1. 1. 要求 OpenCL 实施分配部分内存 在这一示例中，我们需要三个内存块（两个输入，一个输出）。我们使用 C/C++ 实施中的阵列。要分配这些阵列，我们将：

   2.  

       

      1 2 3 4 5 6

      /* Number of elements in the arrays of input and output data. */     int arraySize = 1000000;     /* Arrays to hold the input and output data. */     int* inputA = new int[arraySize];     int* inputB = new int[arraySize];     int* output = new int[arraySize];

      在 OpenCL 中，我们使用内存缓冲区，它们是某一特定大小的内存块。要分配这些缓冲区，我们将：

       

      1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23

      /* Number of elements in the arrays of input and output data. */     cl_int arraySize = 1000000;     /* The buffers are the size of the arrays. */     size_t bufferSize = arraySize * sizeof(cl_int);     /*      * Ask the OpenCL implementation to allocate buffers for the data.      * We ask the OpenCL implemenation to allocate memory rather than allocating      * it on the CPU to avoid having to copy the data later.      * The read/write flags relate to accesses to the memory from within the kernel.      */     bool createMemoryObjectsSuccess = true;     memoryObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);     createMemoryObjectsSuccess &= checkSuccess(errorNumber);     memoryObjects[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);     createMemoryObjectsSuccess &= checkSuccess(errorNumber);     memoryObjects[2] = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);     createMemoryObjectsSuccess &= checkSuccess(errorNumber);     if (!createMemoryObjectsSuccess)     {         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);         cerr << "Failed to create OpenCL buffer. " << __FILE__ << ":"<< __LINE__ << endl;         return 1;     }

      虽然这看上去很复杂，但只有三个 OpenCL API 调用。其差别在于，此处我们将检查是否有错误（这是一种最佳做法），而 C/C++ 中我们不会。

      将内存映射到本地指针 现在内存已经分配好，但只有 OpenCL 实施知道其位置。要访问 CPU 上的缓冲区，我们需要将它们映射到指针上：

       

      1 2 3 4 5 6 7 8 9 10 11 12

      /* Map the memory buffers created by the OpenCL implementation to pointers so we can access them on the CPU. */     bool mapMemoryObjectsSuccess = true;     cl_int* inputA = (cl_int*)clEnqueueMapBuffer(commandQueue, memoryObjects[0], CL_TRUE, CL_MAP_WRITE, 0, bufferSize, 0, NULL, NULL, &errorNumber);     mapMemoryObjectsSuccess &= checkSuccess(errorNumber);     cl_int* inputB = (cl_int*)clEnqueueMapBuffer(commandQueue, memoryObjects[1], CL_TRUE, CL_MAP_WRITE, 0, bufferSize, 0, NULL, NULL, &errorNumber);     mapMemoryObjectsSuccess &= checkSuccess(errorNumber);     if (!mapMemoryObjectsSuccess)     {        cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);        cerr << "Failed to map buffer. " << __FILE__ << ":"<< __LINE__ << endl;        return 1;     }

   3. 这些指针现在可以作为一般的 C/C++ 指针使用。
   4. 在 CPU 上初始化数据 由于我们拥有了指向内存的指针，这一步骤与 CPU 上的相同：
      
       

      1 2 3 4 5

      for (int i = 0; i < arraySize; i++)     {        inputA[i] = i;        inputB[i] = i;     }

       
   5. 取消映射缓冲区要使 OpenCL 设备可以使用这些缓冲区，必须从 CPU 取消它们的映射：
      
       

      1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18

      /*      * Unmap the memory objects as we have finished using them from the CPU side.      * We unmap the memory because otherwise:      * - reads and writes to that memory from inside a kernel on the OpenCL side are undefined.      * - the OpenCL implementation cannot free the memory when it is finished.      */     if (!checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[0], inputA, 0, NULL, NULL)))     {        cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);        cerr << "Unmapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;        return 1;     }     if (!checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[1], inputB, 0, NULL, NULL)))     {        cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);        cerr << "Unmapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;        return 1;     }

       
   6. 将数据映射到内核必须告知内核哪些数据将用作其输入，然后才能调度其运行。此时，我们将内存对象映射到 OpenCL 内核的参数：
      
       

      1 2 3 4 5 6 7 8 9 10

      bool setKernelArgumentsSuccess = true;     setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 0, sizeof(cl_mem), &memoryObjects[0]));     setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 1, sizeof(cl_mem), &memoryObjects[1]));     setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 2, sizeof(cl_mem), &memoryObjects[2]));     if (!setKernelArgumentsSuccess)     {         cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);         cerr << "Failed setting OpenCL kernel arguments. " << __FILE__ << ":"<< __LINE__ << endl;         return 1;     }

       
   7. 运行内核 对于内核代码及其调度方式，请参见基本信息。
   8. 获取结果 当计算完成时，我们按照与映射输入缓冲区相同的方式映射输出缓冲区。然后，我们可以按照一般方式使用指针读取结果，再和之前一样取消映射缓冲区。