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设备计算出结果。

缓冲与子缓冲示意图

Camera课程

Python教程

Java教程

Web教程

数据库教程

图形图像教程

办公软件教程

Linux教程

计算机教程

大数据教程

开发工具教程