OpenCL入门程序

注意：   如果是从显存到显存不是用writebuffer而是用copybuffer

 

 以前就听说OPenCL，今天就特地使用了一下，我的机器是N卡，首先装上了CUDA的开发包，由于CUDA对OPenCL支持比较好，就选择了N卡上的GPU并行计算。

     OPenCL是一个开放的标准和规范，全程是开放计算库，主要是发挥计算机的所有计算资源，包括CPU、GPU、多核等。所以说OPenCL是一个跨硬件和软件平台的开放标准，在此框架下开发的并行计算程序很容易就能移植到其他平台上，也许是这样吧。其实，关于GPU的并行计算的大致思路一般都是CPU向GPU发送一个计算指令，然后把数据拷贝的GPU的显存中参与计算，然后将计算好的显存中的数据拷贝到主机内存中，虽然说，过程大概就是这样，但是其中涉及到的细节可是特别多。下面就以一个简单的例子为例讲述OPenCL编程开发的一般步骤和模型。

    第一步，首先获得可以参与计算的OPenCL平台个数

cl_uint numPlatforms = 0;   //GPU计算平台个数  cl_platform_id platform = NULL;  clGetPlatformIDs(0,NULL,&numPlatforms);

      第二步，获得平台的列表，并选择其中的一个作为计算的平台  

//获得平台列表  cl_platform_id * platforms = (cl_platform_id*)malloc(numPlatforms * sizeof(cl_platform_id));  clGetPlatformIDs (numPlatforms, platforms, NULL);  //轮询各个opencl设备  for (cl_uint i = 0; i < numPlatforms; i ++)  {    char pBuf[100];    clGetPlatformInfo(platforms[i],CL_PLATFORM_NAME,sizeof(pBuf),pBuf,NULL);    printf("%s\n",pBuf);    platform = platforms[i];  }  free(platforms);

第三步，获得硬件设备以及生成上下文  

//获得GPU设备  cl_device_id device;  status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);  //生成上下文  cl_context context = clCreateContext(0, 1, &device, NULL, NULL, &status);

至此，OPenCL的初始化工作已经完成，最好将这个过程封装成一个函数。  

第四步，装载内核程序代码以及生成program

//装载内核程序  size_t szKernelLength = 0;  size_t sourceSize[] = {strlen(kernelSourceCode1)};  char *cFileName = "kernel.cl";  char * cPathAndName= shrFindFilePath(cFileName, argv[0]);  const char* kernelSourceCode = oclLoadProgSource(cPathAndName, "", &szKernelLength);  cl_program program = clCreateProgramWithSource(context,1,&kernelSourceCode,&szKernelLength,&status);  //为所有指定的设备生成CL_program  status = clBuildProgram(program,1,&device,NULL,NULL,NULL);    size_t len = 0;  char buf[2048];  if (status != CL_SUCCESS)  {    status = clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,sizeof(buf),buf,&len);    printf("%s\n", buf);    exit(1);  }

第五步，创建一个命令队列，将这个命令队列放入内核程序中执行  

//创建一个opencl命令队列  cl_command_queue commandQueue = clCreateCommandQueue(context,device,0,&status);  //创建opencl buffer对象  cl_mem outputBuffer = clCreateBuffer(context,CL_MEM_ALLOC_HOST_PTR,4*4*4,NULL,&status);  //得到指定名字的内核实例句柄  cl_kernel kernel = clCreateKernel(program,"hellocl",&status);  //为内核程序设置相应的参数,也就是函数传参  status = clSetKernelArg(kernel,0,sizeof(cl_mem),&outputBuffer);  //将一个kernel放入队列  size_t globalThreads[] = {4,4};  size_t localThreads[] = {2,2};  //开始在设备上执行核函数  status = clEnqueueNDRangeKernel(commandQueue,kernel,2,NULL,    globalThreads,localThreads,0,NULL,NULL);  status = clFinish(commandQueue);

第六步，将计算结果拷贝到主存中  

//将GPU本地内存中的数据拷回到host端的内存中  unsigned int *outbuffer = new unsigned int[4*4];  memset(outbuffer,0,4*4*4);  status = clEnqueueReadBuffer(commandQueue,outputBuffer,    CL_TRUE,0,4*4*4,outbuffer,0,NULL,NULL);

第七步，显示及清理内存  

printf("out:\n");  for (int i = 0; i < 16; i ++)  {    printf("%x ",outbuffer[i]);    if ((i+1)%4 == 0)    {      printf("\n");    }  }  //清理部分  status = clReleaseKernel(kernel);  status = clReleaseProgram(program);  status = clReleaseMemObject(outputBuffer);  status = clReleaseCommandQueue(commandQueue);  status = clReleaseContext(context);  delete outbuffer;

       

核函数如下：

__kernel void hellocl (__global uint *buffer){    uint dim = get_work_dim();//获得工作空间的维度信息  size_t gidx,gidy,gidz;  size_t gsizx,gsizy,gsizz;  if(dim == 1)  {    gidx = get_global_id(0);    gsizx = get_global_size(0);    buffer[gidx] = gidx;  }  else if(dim == 2)  {    gidx = get_global_id(0);    gidy = get_global_id(1);    gsizx = get_global_size(0);    gsizy = get_global_size(1);    buffer[gidx+gidy*gsizx] = (1<<gidx)|(0x10<<gidy);;  }  else   {    gidx = get_global_id(0);    gidy = get_global_id(1);    gidy = get_global_id(2);    gsizx = get_global_size(0);    gsizy = get_global_size(1);    gsizz = get_global_size(2);    buffer[gidx + gidy*gsizx + gidz*gsizx*gsizy] = gidx;  }}

运算结果显示如下图：  

 

这只是最简单的程序，复杂算法的并行化还需要深入研究。