OpenCL 内建管道读/写函数

OpenCL 2.0中引入了一种新的工作机制用于在不同内核间传递数据,这种新的机制就是管道(pipe)。

一个管道实际上就是一个结构化FIFO缓冲,是通过类型修饰符pipe关键字来标识,由数据包(packet)的集合空间构成。这些数据包在管道中是有序放置的,管道只能被内核函数访问,而不能被主机访问。
例如:

pipe int4 pipeA;//int4包的一个管道
pipe user_type_t pipeB;//用户自定包的一个管道

可以使用OpenCL 函数修饰符中的访问修饰符read_only或write_only来限定管道在内核中的读写权限。默认修饰符是read_only。内核不能对同一个管道又读又写,所以管道用read_write修饰符会有编译错误。

管道(即pipe对象)只能作为函数(包含内核函数)参数传入,而不能用在函数内部进行声明或作为程序全局对象使用。

OpenCL 2.0中,增加了管道操作的内建函数,下面我们就分别来讲述这些内建函数。

内建管道读/写函数

下表列出了内建管道读和写函数。在开始讲解这些函数前,我们做一个约定:使用通用类型符gentype来指示这些函数参数可以取OpenCL内建矢量、内建标量、浮点数和用户自定义数据类型。

内建管道读/写函数

下面我们通过一个简单的例子来说明上表中函数的用法。例子中有两个内核函数,分别是pipe_producer和pipe_consumer,pipe_producer向管道中写入数据,而pipe_consumer从管道中读取数据。主机和设备端代码如下:



内核代码:

kernel void pipe_producer(global float *src,
{
}
                              write_only pipe float out_pipe)
    int gid = get_global_id(0);
    reserve_id_t res_id;
    //每个work-item向预留区域管道中读取1个包数据
    res_id = reserve_write_pipe(out_pipe, 1);
    float src_pipe = src[gid];
    if(is_valid_reserve_id(res_id))
    {
        //把src_pipe数据写入管道中
        if(write_pipe(out_pipe, res_id, 0, &src_pipe) != 0)
              return;
        //提交写入操作
        commit_write_pipe(out_pipe, res_id);
    }
}
kernel void pipe_consumer(global float *dst,
{
read_only pipe float in_pipe)
 int gid = get_global_id(0);
 reserve_id_t res_id;
 //每个work-item从预留区域管道中写入1个包数据
 res_id = reserve_read_pipe(in_pipe, 1);
 float dst_pipe;
 if(is_valid_reserve_id(res_id))
 {
   //从管道中读取1个包数据到dst_pipe
   if(read_pipe(in_pipe, res_id, 0, &dst_pipe) != 0)
      return;
  //提交读取操作
  commit_read_pipe(in_pipe, res_id);
 }
 dst[gid] = dst_pipe;
}

主机端代码:

……
int numPackets = 16384;
int packetSize = sizeof(float);
kernel_pipe_consumer = clCreateKernel(program, "pipe_consumer",
&err);
checkErr(err, _LINE_) ;
kernel_pipe_producer = clCreateKernel(program, "pipe_producer",
                                              &err);
checkErr(err, _LINE_) ;
//细粒度SVM
src = (float *)clSVMAlloc(context, CL_MEM_READ_WRITE |
                              CL_MEM_SVM_FINE_GRAIN_BUFFER, sizeBytes,
                              0);
dst = (float *)clSVMAlloc(context, CL_MEM_READ_WRITE |
                              CL_MEM_SVM_FINE_GRAIN_BUFFER, sizeBytes,
                              0);
if(src == NULL || dst == NULL)
{
    printf("allocate SVM error!\n");
    return 0;
}
//创建管道,包数据类型为float,包个数为numPackets个
pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, packetSize,
                        numPackets, NULL, &err);
checkErr(err, _LINE_);
for (int i = 0; i 〈 numPackets; i++)
{
    src[i] = i;
    dst[i] = 0.0f;
}
size_t globalworksize = numPackets;
size_t localworksize = 128;
//producer
err = clSetKernelArgSVMPointer(kernel_pipe_producer, 0, src);
checkErr(err, _LINE_);
err = clSetKernelArg(kernel_pipe_producer, 1, sizeof(cl_mem),
                        &pipe);
checkErr(err, _LINE_);
err = clEnqueueNDRangeKernel(cmdQueue0, kernel_pipe_producer, 1,
                                  NULL,
                                  &globalworksize, &localworksize, 0,
                                  NULL, &producer_event);
checkErr(err, _LINE_);
//consumer
err = clSetKernelArgSVMPointer(kernel_pipe_consumer, 0, dst);
checkErr(err, _LINE_);
err = clSetKernelArg(kernel_pipe_consumer, 1, sizeof(cl_mem),
                        &pipe);
checkErr(err, _LINE_);
err = clEnqueueNDRangeKernel(cmdQueue1, kernel_pipe_consumer, 1,
                                  NULL,
&globalworksize, &localworksize, 1,
&producer_event, NULL);
checkErr(err, _LINE_);
err = clFinish(cmdQueue1);
checkErr(err, _LINE_);

上述代码使用管道实现了生产-消费模型,管道作为参数传入pipe_producer内核,内核中产生数据写入管道中;在pipe_consumer内核中,从管道读取pipe_producer内核中生成的写入的数据。在pipe_producer中每个工作项向管道中写入1个包数据,需要注意的是写入管道中的顺序是乱序的。OpenCL执行模型中,工作组的执行是相互独立的,我们无法控制全局工作项间读取/写入管道的顺序,但是我们可以在工作组内控制读取/写入次序,具体内核代码如下:

kernel void pipe_producer(global float *src,
{
}
                              write_only pipe float out_pipe)
    int gid = get_global_id(0);
    local reserve_id_t res_id;
    if(get_local_id(0) == 0)
        //预留区域管道写入get_local_size(0)个包
        res_id = reserve_write_pipe(out_pipe, get_local_size(0));
    barrier(CLK_LOCAL_MEM_FENCE);
    float src_pipe = src[gid] ;
    if(is_valid_reserve_id(res_id))
    {
        //根据工作项工作组内索引id写入管道对应位置
        if(write_pipe(out_pipe, res_id, get_local_id(0), &src_pipe)
            != 0)
              return;
        commit_write_pipe(out_pipe, res_id);
    }
kernel void pipe_consumer(global float *dst,
{
                              read_only pipe float in_pipe)
    int gid = get_global_id(0);
    local reserve_id_t res_id;
    if(get_local_id(0) == 0)
        //预留区域管道读取get_local_size(0)个包
        res_id = reserve_read_pipe(in_pipe, get_local_size(0));
    barrier(CLK_LOCAL_MEM_FENCE);
    float dst_pipe;
    if(is_valid_reserve_id(res_id))
    {
      //根据工作项工作组内索引id读取管道对应位置
      if(read_pipe(in_pipe, res_id, get_local_id(0), &dst_pipe) != 0)
        return;
        commit_read_pipe(in_pipe, res_id);
}
dst[gid] = dst_pipe;
}

在上述代码中,我们使用局部存储器声明预留区域ID res_id,在每个工作组中的第一个工作项确定预留区域写入get_local_size(0)个包,在写入和读取管道时根据工作项在工作组内的索引来确定管道对应位置,从而达到了对工作组内管道数据的顺序写入和读取。
除了使用局部存储器,我们也可以使用工作组管道读/写函数来实现工作组内的顺序操作。

赞(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映射