GPU的OPENCL编程2(CPU和GPU数据交互和GPU与global id local id的关系)

在不同的架构上GPU和CPU的情况不一样，在电脑上GPU可以有很大的DDR显存存放数据,SOC上GPU的显存和CPU是同一个DDR，GPU可能会提供一个小的SRAM作为缓存该缓存不用于CPU和GPU交互数据。这个章节的细节内容建议查看“OpenCL Programming Guide”这本书。
在Opencl编程中和普通CPU编程有很大的不同，提供给GPU的共享内存需要通过clCreateBuffer或clCreateImage申请。传输变量内存地址，传输Kernel的参数都必须通过clSetKernelArg函数调用。CPU和GPU间的数据共享，参数传递都必须通过封装好的函数处理，不能直接简单的用变量名称共享。

图1 ：CPU和GPU的内存关系

CPU与GPU的内存空间交互

申请给CPU申请和GPU交互的DDR内存空间主要是以下3种方式：
1.clCreateBuffer函数是创建缓存对象是与图像处理无关内存。
2. clCreateImage2D函数是创建2D图像格式空间内存。
3. clCreateImage3D函数是创建3D图像格式空间内存。
对于Image的创建的内存使用一些和Image相关的函数调用比如：clEnqueueMapImage、clEnqueueReadImage、SaveImage、可以参考“Images and Samplers章节”。
对于为什么需要Image这样的一个内存对象我认为主要是方便使用，对于图像 Specifying and querying for image formats（可以设置和查询图像格式），可以使用sampler快速调用函数获取你所需要的图像,而无需像内存图像一样各种计算后你才能获取到图像所在的位置指针。clGetMemObjectInfo能防护图像信息又能返回内存对象信息，可以将获取图像后方便转换为内存操作。

CPU对申请的Buffer和Image空间访问可以通过2种方式。
1.通过Buffer和Image对应的read*和write函数进行访问。
2.通过Buffer和Image对应的**map内存映射函数进行内存直接访问。

内存数据划分

使用Opencl的目的就是为了并行计算，大量的数据并行计算一个整块的数据如何分配给每个GPU对并行的性能提升非常重要。GPU的优势为重复执行相同的算法操作，判断语句和跳转语句在GPU上运行效率奇低。我们做一个大数组排序、取最大值等运算时会用到内存划分。数组不用内存划分的话可能需要使用FOR循环的方式很难发挥GPU并行运算的效能。clEnqueueNDRangeKernel是clEnqueueKernel执行内核控制的扩展。查看“Writing a Data-Parallel Kernel Using OpenCL C”章节。
一个小的程序例子进行一下说明有帮助理解：
C语言的单个循环例子：

void scalar_add (int n, const float *a, const float *b, float *result)
{
 int i;
for (i=0; i<n; i++)
 result[i] = a[i] + b[i];
}

在使用clEnqueueNDRangeKernel后可以变成

kernel void
scalar_add (global const float *a, 
global const float *b, 
global float *result)
{
 int id = get_global_id(0);
 result[id] = a[id] + b[id];
}

C语言的2层个循环例子：

void scalar_add (int n, const float *a, const float *b, float *result)
{
 int i，j;
 for(j=0;j<n;j++)
	for (i=0; i<n; i++)
	 result[i][j] = a[i][j] + b[i][j];
}

在使用clEnqueueNDRangeKernel后可以变成

kernel void
scalar_add (global const float *a, 
global const float *b, 
global float *result)
{
 int id = get_global_id(0);
 int id1 = get_global_id(1);
 result[id][j] = a[id][j] + b[id][j];
}

从上面可以看到kernel消除了for循环，使用了get_global_id 获取计算时候的维度信息。期中相关维度信息是在clEnqueueNDRangeKernel的参数中设置的。

cl_int clEnqueueNDRangeKernel (	cl_command_queue command_queue,
 	cl_kernel kernel,
 	cl_uint work_dim,
 	const size_t *global_work_offset,
 	const size_t *global_work_size,
 	const size_t *local_work_size,
 	cl_uint num_events_in_wait_list,
 	const cl_event *event_wait_list,
 	cl_event *event)

对应的设置如下图一个例子理解：
给出了内核如何访问 clEnqueueNDRangeKernel 中指定的全局和本地工作大小的示例在设备上。在此示例中，内核在全局工作大小上执行16 个项目，工作组大小为每组 8 个项目。OpenCL 没有描述全局和本地 ID 如何映射到工作项和工作组。

图2 clEnqueueNDRangeKernel 对应的work-item

clEnqueueNDRangeKernel 中对于我使用的SOC，global_work_size和local_work_size相除受group size 设置收到了很大的限制，设置异常会否则运行会报错-54 CL_INVALID_WORK_GPROUP_SIZE。
我使用的SOC的 local_work_size最大为128。
所以建议在编写程序时使用clGetKernelWorkGroupInfo查询对应CL_KERNEL_WORK_GROUP_SIZE、CL_KERNEL_LOCAL_MEM_SIZE、CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_
MULTIPLE、CL_KERNEL_PRIVATE_MEM_SIZE。以免编程时出现设置变量、组超界的问题。

对于global id 和local id 和计算的关系

这个就表示上面这个数据并行计算的kernel中每一个work-group由8个work-item组成（看图2理解），而共有16个work-items要被处理，即总的work-items要被分成2个work-group。另外work-item对应硬件上的一个PE（processing element）,而一个work-group对应硬件上的一个CU（computing unit）。这种对应可以理解为，一个work-item不能被拆分到多个PE上处理；同样，一个work-group也不能拆分到多个CU上同时处理。当映射到OpenCL硬件模型上时，每一个work-item运行在一个被称为处理基元（processing element）的抽象硬件单元上，其中每个处理基元可以处理多个work-item(注：摘自《OpenCL异构计算》P87)。

对应前面的图使用一个OPENCL程序能很好的说明并理解global id 和local id 的关系：

  int global_id_0=get_global_id(0);\
  int local_id_0=get_local_id(0);\
  int global_size_0=get_global_size(0);\
  int local_size_0=get_local_size(0);\
  int write_addr=global_id_0;\
  float data_id=(float)local_id_0*0.1+(float)(global_id_0);\
  pdst[write_addr]=data_id;\

对数组输出的结果
整数部分为global_size_0 ，小数部分为0.1位是local_id_0。
从结果可以看出get_global_id=0的是GPU0处理 ，get_global_id=1的是GPU1处理类推，到get_global_id=8时又个CPU0处理，以此类推循环。

fdst[0]=0.000000
fdst[1]=1.100000
fdst[2]=2.200000
fdst[3]=3.300000
fdst[4]=4.400000
fdst[5]=5.500000
fdst[6]=6.600000
fdst[7]=7.700000
fdst[8]=8.000000
fdst[9]=9.100000
fdst[10]=10.200000
fdst[11]=11.300000
fdst[12]=12.400000
fdst[13]=13.500000
fdst[14]=14.600000
fdst[15]=15.700000