3.8 CUDA编程者使用OpenCL的注意事项

英伟达的CUDA C提供的API与OpenCL类似。代码清单3.6用使用CUDA实现了向量相加，OpenCL和CUDA中的很多命令都可以相互对应。OpenCL API中有更多的参数，这是因为OpenCL需要在运行时去查找平台，并对程序进行编译。CUDA C只针对英伟达的GPU，其只有一个平台可以使用，所以平台的查找自动完成；将程序编译成PTX的过程可以在编译主机端二进制文件的时候进行(CUDA没有运行时编译)。

OpenCL中平台在运行时进行查找，程序需要选择一个目标设备，并在运行时进行编译。程序的编译不能在运行时之外完成，因为不知道哪个具体设备要去执行内核，无法在这种情况下生成中间码(IL/ISA)。比如，一个OpenCL内核在AMD GPU上测试有没有问题，但当其需要运行在Intel的设备上时，编译器生成的中间码就和AMD平台上的不太一样。从而，对与平台的查找，以及在运行时进行编译程序就能避免这样的问题。

OpenCL和CUDA C最大的区别是，CUDA C提供一些特殊的操作完成内核启动，其编译需要使用一套工具链，这套工具链中需要包含英伟达支持的预处理器。预处理器生成的代码，和OpenCL的内核代码十分相近。

  1. #include <stdio.h>
  2. #include <stdlib.h>
  3. #include <math.h>
  5. // CUDA kernel. Eache thread computes one element of C
  6. __global__ void vecAdd(int *A, int *B, int *C, int elements){
  7.   // Compute the global thread ID using CUDA intrinsics
  8.   int id = blocakIdx.x * blocakDim.x + threadIdx.x;
  10.   // Must check that the thread is not out of bounds
  11.   if (id < elements)
  12.     C[id] = A[id] + B[id];
  13. }
  15. int main(int argc, char *argv[]){
  16.   // Elements in each array
  17.   const int elements = 2048;
  19.   // Compute the size of the data
  20.   size_t datasize = sizeof(int) * elements;
  22.   // Allocate space for input/output host data
  23.   int *A = (int *)malloc(datasize); // Input array
  24.   int *B = (int *)malloc(datasize); // Input array
  25.   int *C = (int *)malloc(datasize); // Output array
  27.   // Device input vectors
  28.   int *bufA;
  29.   int *bufB;
  30.   // Device output vectors
  31.   int *bufC;
  33.   // Allocate memeory for each vector on GPU
  34.   cudaMalloc(&bufA, datasize);
  35.   cudaMalloc(&bufB, datasize);
  36.   cudaMalloc(&bufC, datasize);
  38.   int i;
  39.   // Initialize vectors on host
  40.   for (i = 0; i < elements; i++){
  41.     A[i] = i;
  42.     B[i] = i;
  43.   }
  45.   // Copy host vectors to device
  46.   cudaMemcpy(bufA, A, datasize, cudaMemcpyHostToDevice);
  47.   cudaMemcpy(bufB, B, datasize, cudaMemcpyHostToDevice);
  49.   int blockSize, gridSize;
  51.   // Number of threads in each thread block
  52.   blockSize = 256;
  54.   // Number of thread blocks in grid
  55.   gridSize = elements / blockSize;
  57.   // Execute the kernel
  58.   vecAdd<<<gridSize, blockSize>>>(bufA, bufB, bufC, elements);
  60.   // Copy array back to host
  61.   cudaMemcpy(C, bufC, datasize, cudaMemcpyDeviceToHost);
  63.   // Release device memeory
  64.   cudaFree(bufA);
  65.   cudaFree(bufB);
  66.   cudaFree(bufC);
  68.   // Release host memeoy
  69.   free(A);
  70.   free(B);
  71.   free(C);
  72. }

代码清单3.6 CUDA C版本的向量相加