OpenCL 创建子缓冲区对象|极客笔记

针对在OpenCL 分配缓冲区对象中创建的缓冲区对象，我们可以进一步划分为子缓冲区，在不同的设备上分配子缓冲区。利用如下函数创建子缓冲区：

cl_mem clCreateSubBuffer (cl_mem buffer,
                              cl_mem_flags flags,
                              cl_buffer_create_type buffer_create_type,
                              const void ＊buffer_create_info,
                              cl_int ＊errcode_ret)

参数buffer_create_type,buffer_create_info共同描述所要创建的缓冲区对象的类型。buffer_create_type的取值由枚举cl_buffer_create_type定义，其值类型是CL_BUFFER_CREATE_TYPE_REGION。

创建一个缓冲区对象，表示buffer中的一个特定区域。

buffer_create_info是如下结构体的指针：

typedef struct _cl_buffer_region{
  size_t origin;
  size_t size;
}cl_buffer_region

cl_buffer_region中的{origin,size}使用字节数定义了buffer中的偏移量和大小。如果使用CL_MEM_USE_HOST_PTR创建buffer，那么与所返回的缓冲区对象关联的host_ptr为host_ptr+origin。返回的缓冲区对象会引用为buffer分配的数据，并指向这个数据中（origin,size）指定的一个特性区域。

在结构体cl_buffer_region中，size定义了子缓冲区的大小，origin指定了子缓冲在缓冲区中的开始位置。需要注意的是，origin指定的开始位置必须对齐到clGetDeviceInfo（）中的CL_DEVICE_MEM_BASE_ADDR_ALIGN值，否则返回的错误码为CL_MISALIGNED_SUB_BUFFER_OFFSET。也就是说，如果子缓冲设备的开始位置不满足对齐，则无法使用子缓冲来划分缓冲区域。

在支持OpenCL的设备有AMD FX-8320和AMD R9 285 GPU。在如下代码中，笔者把创建的缓冲划分为两个子缓冲，用两个不同的OpenCL设备去处理子缓冲区的数据，代码如下：

内核代码：

//simple.cl
__kernel void MultAdd(__global ＊buffer)
{
    size_t id = get_global_id(0);
    buffer[id] = buffer[id] + buffer[id];
}

主机代码：

……
cl_platform_id platform;
cl_device_id ＊devices;
cl_uint numDevice;
cl_context context;
cl_command_queue ＊cmdQueue;
cl_int err;
cl_program program;
cl_kernel kernel;
cl_event event[2];
err = clGetPlatformIDs(1, &platform, NULL);
//获得平台上所有的OpenCL设备,第一个设备为GPU,第二个设备为CPU
err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL,
&numDevice);
devices = (cl_device_id ＊)malloc(sizeof(cl_device_id) ＊
                                        numDevice);
err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL,
                        numDevice, devices, NULL);
context = clCreateContext(NULL, numDevice, devices, NULL, NULL,
                              &err);
//每个设备处理的数据个数
int NumBufferDevice = 512;
int DataSize = NumBufferDevice ＊ numDevice;
int ＊d_SrcData;
d_SrcData = (int ＊)malloc(sizeof(int) ＊ DataSize);
for(int i = 0; i 〈 DataSize; i++)
{
    d_SrcData[i] = i;
}
//创建缓冲区
cl_mem SrcData;
SrcData = clCreateBuffer(context,
                              CL_MEM_READ_WRITE |
                              CL_MEM_COPY_HOST_PTR,
                              sizeof(int) ＊ DataSize, d_SrcData, &err);
cl_mem ＊subuffer;
subuffer = (cl_mem ＊)malloc(sizeof(cl_mem) ＊ numDevice);
cmdQueue = (cl_command_queue ＊)malloc(
                  sizeof(cl_command_queue) ＊ numDevice);
for(int i = 0; i 〈 numDevice; i++)
{
    //为每个设备创建命令队列
cmdQueue[i] = clCreateCommandQueue(context, devices[i], NULL,
                                          &err);
    //创建子缓冲区
    cl_buffer_region region;
    region.origin = NumBufferDevice ＊ i ＊ sizeof(int);
    region.size = NumBufferDevice ＊ sizeof(int);
    subuffer[i] = clCreateSubBuffer(SrcData,
                                          CL_MEM_READ_WRITE,
                                          CL_BUFFER_CREATE_TYPE_REGION,
                                          &region, &err);
}
size_t program_length;
char ＊const source = readKernelFile("simple.cl",
                                          &program_length);
program = clCreateProgramWithSource(context, 1,
                                          (const char ＊＊)&source,
                                          NULL, &err);
err = clBuildProgram(program, numDevice, devices, NULL, NULL,
                        NULL);
kernel = clCreateKernel(program, "MultAdd", &err);
for(int i = 0; i 〈 numDevice; i++)
{
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &subuffer[i]);
    size_t globalsize[] = {NumBufferDevice};
    err = clEnqueueNDRangeKernel(cmdQueue[i], kernel, 1,
                                        NULL, globalsize, NULL,
                                        0, NULL, &event[i]);
}
//拷贝缓冲数据
err = clEnqueueReadBuffer(cmdQueue[0], SrcData, CL_TRUE,
                              0, sizeof(int) ＊ DataSize, d_SrcData,
                              2, event, NULL);
//验证结果
for(int i = 0; i 〈 DataSize; i++)
{
    if(d_SrcData[i] != 2 ＊ i)
{
    printf("error at %d\n", i);
break;
}
}

上述代码中，子缓冲与缓冲的关系可以用下图表示。在缓冲区中有512×2个数据，包含前512个数据的子缓冲属于GPU，包含后512个数据的子缓冲属于CPU。分别把子缓冲提交给不同的OpenCL设备计算出结果。

缓冲与子缓冲示意图