OpenCL 内建工作组管道读/写函数

除了工作项对管道读/写函数，OpenCL也支持工作组级别的的读/写函数。一个工作组内所有工作项必须都执行工作项管道读/写函数，并且这些函数的参数值相同。下表中列出了这些工作组管道读/写函数。在开始讲解这些函数前，我们做一个约定：使用通用类型符gentype来指示这些函数参数可以取OpenCL内建矢量、内建标量、浮点数和用户自定义数据类型。

还是以OpenCL 内建管道读/写函数的例子为例，为了实现在工作组内对管道的顺序操作，我们使用上表中的内建函数，具体代码如下：

kernel void pipe_producer(global float ＊src,
                              write_only pipe float out_pipe)
{
    int gid = get_global_id(0);
    reserve_id_t res_id;
    //保留区域内写入get_local_size(0)个包
    res_id = work_group_reserve_write_pipe(out_pipe,
                                                    get_local_size(0));
    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;
      work_group_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;
  //保留区域内写入get_local_size(0)个包
  res_id = work_group_reserve_read_pipe(in_pipe,
                                        get_local_size(0));
  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;
     work_group_commit_read_pipe(in_pipe, res_id);
  }
  dst[gid] = dst_pipe;
}

使用上表中的内建工作组管道读/写函数，相比较使用全局存储器实现方案，上述代码更简洁。