OpenCL 矢量存取

▶ 函数 vloadn 和 vstoren 来实现全局存储器和局部存储器之间的向量拷贝

● 代码

 #include <stdio.h>

 #include <stdlib.h>

 #include <cl.h>

 const int nElement = ;

 const char *programSource = "                           \

 __kernel void prog(__global int *A, __global int *B)    \

 {                                                       \

     int idx = get_global_id();                         \

     int4 temp = vload4(idx, A);                         \

     vstore4(temp, idx, B);                              \

     return;                                             \

 }                                                       \

 ";

 int main()

 {

     const size_t datasize = sizeof(int) * nElement;

     int i, *A, *B;

     cl_int status;

     A = (int*)malloc(datasize);

     B = (int*)malloc(datasize);

     for (i = ; i < nElement; A[i] = i, B[i] = , i++);

     cl_platform_id platform;

     clGetPlatformIDs(, &platform, NULL);

     cl_device_id device;

     clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, , &device, NULL);

     cl_context context = clCreateContext(NULL, , &device, NULL, NULL, &status);

     cl_command_queue cmdQueue = clCreateCommandQueue(context, device, , &status);

     cl_mem bufferA, bufferB;

     bufferA = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status);

     bufferB = clCreateBuffer(context, CL_MEM_WRITE_ONLY, datasize, NULL, &status);

     clEnqueueWriteBuffer(cmdQueue, bufferA, CL_FALSE, , datasize, A, , NULL, NULL);

     cl_program program = clCreateProgramWithSource(context, , &programSource, NULL, &status);

     status = clBuildProgram(program, , &device, NULL, NULL, NULL);

     cl_kernel kernel = clCreateKernel(program, "prog", &status);

     clSetKernelArg(kernel, , sizeof(cl_mem), &bufferA);

     clSetKernelArg(kernel, , sizeof(cl_mem), &bufferB);

     size_t globalSize[] = { nElement }, localSize[] = {  };

     status = clEnqueueNDRangeKernel(cmdQueue, kernel, , NULL, globalSize, localSize, , NULL, NULL);

     clEnqueueReadBuffer(cmdQueue, bufferB, CL_TRUE, , datasize, B, , NULL, NULL);

     for (i = ; i < nElement; i++)

     {

         if (B[i] != i)

             break;

     }

     printf("Output is %s.\n", (i == nElement) ? "correct" : "incorrect");

     free(A);

     free(B);

     clReleaseContext(context);

     clReleaseMemObject(bufferA);

     clReleaseMemObject(bufferB);

     clReleaseCommandQueue(cmdQueue);

     clReleaseProgram(program);

     clReleaseKernel(kernel);

     getchar();

     return ;

 }

● 输出结果

Output is correct.

● 教训

■ 核函数代码中每个 "\" 的后面不要有任何东西，包括空格。因为 "\" 在预处理以后会消失，其后的内容会被当成下一行的内容，而空格会在IDE中使 "\" 失效，导致编译错误
■ 核函数代码中不要有 "//" 型的行注释，理由类似。会使得 "//" 以后的代码全部失效
■ 可以改用字符串连接来写核函数代码，如：

 const char *programSource =

 "__kernel void prog(__global int *A, __global int *B)   \n"

 "{                                                      \n"

 "    int idx = get_global_id(0);                        \n"

 "    int4 temp = vload4(idx, A);                        \n"

 "    vstore4(temp, idx, B);                             \n"

 "    return;                                            \n"

 "}                                                      \n"

 ";                                                      \n"