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;
  }
}

运算结果显示如下图：  

 

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