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

缓冲与子缓冲示意图

赞(1)
未经允许不得转载:极客笔记 » OpenCL 创建子缓冲区对象

评论 抢沙发

  • 昵称 (必填)
  • 邮箱 (必填)
  • 网址
OpenCL 基本概念
OpenCL 是什么OpenCL 平台模型OpenCL 执行模型OpenCL 上下文简介OpenCL 命令队列简介OpenCL 在设备上执行内核OpenCL 存储器区域OpenCL 存储器对象OpenCL 共享虚拟存储器OpenCL 与OpenGL
OpenCL 基础教程
OpenCL 在Windows上搭建开发环境OpenCL 在Linux上搭建开发环境OpenCL 在OS X上搭建开发环境OpenCL 第一个程序OpenCL 平台OpenCL 设备OpenCL 创建上下文OpenCL 创建命令队列OpenCL 创建Program对象OpenCL 编译Program对象OpenCL 查询和管理Program对象OpenCL 创建内核对象OpenCL 设置内核参数OpenCL 查询和管理内核对象OpenCL 执行内核OpenCL 编写内核代码OpenCL 错误处理
OpenCL C特性
OpenCL 地址空间修饰符OpenCL 函数修饰符OpenCL 对象访问修饰符OpenCL 标量数据类型OpenCL 为什么要有矢量数据类型OpenCL 矢量初始化OpenCL 读取和修改矢量分量OpenCL 运算符OpenCL 维度和工作项OpenCL 工作组OpenCL 矢量数据拷贝OpenCL 异步拷贝和预取OpenCL 数学函数OpenCL 公共函数OpenCL 几何函数OpenCL 整数函数OpenCL 关系函数OpenCL 杂项矢量函数OpenCL 同步函数OpenCL 原子函数OpenCL 内建图像读函数OpenCL 内建无采样器图像读函数OpenCL 内建图像写函数OpenCL 内建图像查询函数OpenCL 工作组函数OpenCL 内建管道读/写函数OpenCL 内建工作组管道读/写函数OpenCL 内建管道查询函数OpenCL 设备队列OpenCL Blocks语法OpenCL 设备队列相关函数OpenCL 子内核存储器可见性OpenCL 设备队列的使用示例
OpenCL 存储器对象
OpenCL 存储器对象OpenCL 分配缓冲区对象OpenCL 创建子缓冲区对象OpenCL 图像对象和采样器对象OpenCL 图像对象OpenCL 图像格式描述符OpenCL 图像描述符OpenCL 图像对象查询OpenCL 采样器对象OpenCL 主机端采样器对象OpenCL 设备端采样器对象OpenCL 图像旋转示例OpenCL 管道OpenCL 创建管道对象OpenCL 管道对象查询OpenCL 主机与设备间数据传输OpenCL 图像对象主机与设备间数据拷贝OpenCL 缓冲区对象数据填充OpenCL 图像对象数据填充OpenCL 缓冲区对象间数据传输OpenCL 图像对象和缓冲区对象间数据拷贝OpenCL 缓冲区对象映射OpenCL 图像对象映射OpenCL 解映射OpenCL 共享虚拟存储器OpenCL SVM缓冲创建与释放OpenCL SVM缓冲映射与解映射OpenCL SVM缓冲填充与拷贝OpenCL SVM类型OpenCL SVM特性OpenCL 共享虚拟存储器示例OpenCL 存储器一致性模型OpenCL 存储器次序规则OpenCL 原子操作的存储器次序规则OpenCL 栅栏操作的存储器次序规则OpenCL 工作组函数的存储器次序规则OpenCL 主机端与设备端命令的存储器次序规则OpenCL 关于存储器次序在实际OpenCL计算设备中的实现
OpenCL 同步及事件机制
OpenCL 同步及事件机制OpenCL 主机端的OpenCL同步OpenCL OpenCL事件机制OpenCL 对OpenCL事件的标记和栅栏OpenCL 内核程序中的同步OpenCL 工作组内同步OpenCL 原子操作OpenCL 1.2中的原子操作OpenCL 2.0中的原子操作OpenCL 局部存储器与全局存储器间的异步拷贝OpenCL 工作组间同步
OpenCL 与OpenGL互操作
OpenCL 与OpenGL互操作OpenCL 从一个OpenGL上下文来创建OpenCL上下文OpenCL 使用OpenGL共享的缓存对象OpenCL 使用OpenGL纹理数据OpenCL 共享OpenGL渲染缓存OpenCL 从一个OpenCL存储器对象查询OpenGL对象信息OpenCL 访问共享对象的OpenCL与OpenGL之间的同步OpenCL AMD Cayman架构GPUOpenCL AMD GCN架构的GPUOpenCL NVIDIA CUDA兼容的GPUOpenCL NVIDIA GPU架构的执行模型OpenCL NVIDIA GPU的全局存储器OpenCL NVIDIA GPU的局部存储器OpenCL ARM Mali GPU硬件架构OpenCL ARM Mali GPU存储器层次OpenCL ARM Mali GPU OpenCL映射