OpenCL 局部存储器与全局存储器间的异步拷贝

介绍完了原子操作之后,本节我们再来介绍一下OpenCL从全局存储区域到局部存储区域以及从局部存储区域到全局存储区域的异步拷贝。这些操作在OpenCL 1.2中就已经被引入了。我们之前提到了OpenCL主机端的API提供了事件等待的函数接口,而对于大部分clEnqueue系API而言都会带有event_wait_list参数,用于指定等待哪些事件完成才能执行,以及event参数用来跟踪当前命令本身。在主机端,事件对象类型为cl_event。而OpenCL内核程序中也含有事件对象类型,被定义为event_t。事件对象可以用来跟踪异步拷贝的完成情况。
下面先介绍异步工作组基本拷贝功能的函数原型:

event_t async_work_group_copy (__local gentype *dst,
                                    const __global gentype *src, size_t
                                    num_gentypes,
                                    event_t event)
event_t async_work_group_copy (__global gentype *dst,
                                    const __local gentype *src, size_t
                                    num_gentypes,
                                    event_t event)

上一个函数是将全局存储器的数据拷贝到局部存储器中,下一个则是将局部存储器的数据拷贝到全局存储器中。这里,gentype是一个泛型类型,能支持OpenCL中所有标量及向量形式的基本类型。参数num_gentypes指明了需要拷贝多少个gentype的元素,因此一共需要拷贝的字节数为num_gentypes * sizeof(num_gentypes)。最后一个参数event是用于指定之前进行异步拷贝的事件对象。前后相继的异步拷贝操作可以共享一个事件对象。event参数被指定为前某一个事件对象时,该函数将直接返回该event事件对象;否则这个参数传0即可,函数返回对应该操作的一个事件对象。
这里需要注意的是,当我们调用异步拷贝函数时需要当前工作组的所有工作项参与进行操作,否则结果是未定义的。这意味着如果我们用以下代码来调用异步操作,结果是不确定的:

if(get_local_id(0) == 0)
    async_work_group_copy(pDst, tmpBuffer, 64, 0);

然后,我们再介绍一下事件等待的内建函数,其函数原型如下:

void wait_group_events (int num_events, event_t *event_list)

这个函数非常简单,第一个参数num_events指定了第二个参数event_list中包含了多少个事件对象。第二个参数event_list就是指向事件对象的数组。这个函数与主机端的事件等待类似,只有当事件列表中的事件全都处于完成状态之后程序才能继续往下执行。当然,这个函数与async_work_group_copy一样,必须要当前工作组中所有工作项参与进来操作。另外,我们要注意的是,OpenCL标准中没有规定这个函数被调用之后,所有工作项都会处于同一个调用点,因此这个函数不能被当作一个barrier来使用。不过,这个函数可以保证当前工作项的拷贝工作能够切实完成。
下面,我们来举一个简单的例子来看一下基本功能的异步拷贝操作是如何具体使用的。我们将上述主机端代码从ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,……{evt1,evt2},NULL);这一行一直到FINISH标签,替换为如下代码:

ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
                                  (const size_t[]){
contentLength / sizeof(int)},
                                  (const size_t[]){
                                  maxWorkGroupSize}, 2,
                                  (const cl_event[]){
                                  evt1, evt2}, NULL);
if(ret != CL_SUCCESS)
{
    puts("kernel1 execution failed");
}
    goto FINISH;
//这里用clFinish做命令执行同步
clFinish(command_queue);
//准备做校验
pDeviceBuffer = malloc(contentLength);
//这里使用阻塞的方式读取数据
clEnqueueReadBuffer(command_queue, dstMemObj, CL_TRUE, 0,
                      contentLength,
                      pDeviceBuffer, 0, NULL, NULL);
//做数据校验
bool isOK = true;
for(int i = 0; i 〈 contentLength / sizeof(int); i++)
{
}
    if(pHostBuffer[i] * 2 != pDeviceBuffer[i])
    {
        isOK = false;
        break;
  }
puts(isOK ? "OK" : "NG");

主机端代码比较简单,不再赘述。我们马上看一下对应的OpenCL内核程序代码:

__kernel void kernel_test(__global int *pDst,
{
                                    __global int *pSrc1, __global int *pSrc2)
    local int tmpBuffer[GROUP_NUMBER_OF_WORKITEMS];
    const size_t address_offset = get_group_id(0)
                                            * GROUP_NUMBER_OF_WORKITEMS;
    //这里使用了异步拷贝操作,
    //将全局存储器的数据拷贝到相应的当前工作组的局部存储器中
    //这里,函数原型中的gentype就是int类型
    //一共拷贝GROUP_NUMBER_OF_WORKITEMS个元素
    event_t event = async_work_group_copy(tmpBuffer,
                                                      &pSrc1[address_offset],
                                                      GROUP_NUMBER_OF_WORKITEMS, 0);
    //这里抽空可以做其他事情,比如确定当前工作组中当前工作项的id等
    const int index = get_local_id(0);
    //等待拷贝结束
    wait_group_events(1, &event);
    //我们将对应元素乘以2,再写回局部存储器
    tmpBuffer[index] *= 2;
    //我们确保所有工作项都在调用后续的异步操作之前能够都完成对数据的修改操作
    barrier(CLK_LOCAL_MEM_FENCE);
    //将当前局部存储器中的数据再拷贝回对应的全局存储器中
    event = async_work_group_copy(&pDst[address_offset], tmpBuffer,
                                                      GROUP_NUMBER_OF_WORKITEMS, 0);
    //等待拷贝结束
    wait_group_events(1, &event);
}

这部分代码非常简单。尽管我们可以直接用计算-修改的方式去做,不过作为一个简单的demo而言能作为使大家更容易理解的函数使用方式。这里再强调一下,async_work_group_copy作用于整个工作组,而不是某一工作项。因此在调用这个函数时,我们要清楚与其相关的地址偏移以及拷贝的数据大小需要参考的是工作组的区域范围而不是当前工作项的区域范围。
下面我们再介绍更具灵活性的,带有跨度的异步拷贝操作。下面先给出函数原型:

event_t async_work_group_strided_copy(__local gentype *dst,
                                              const __global gentype *src,
                                              size_t num_gentypes, size_t
                                              src_stride,
                                              event_t event)
event_t async_work_group_strided_copy(__global gentype *dst,
                                              const __local gentype *src,
                                              size_t num_gentypes, size_t
                                              dst_stride,
                                              event_t event)

这两个异步拷贝函数原型与上面两个基本差不多,不过这里分别多了一个参数。从全局存储器拷贝到局部存储器的函数中新增了src_stride参数,表示从全局存储器取数据时,一个元素与其后面的一个元素之间跨多少元素。例如,如果gentype是f loat类型,并且src_stride是4,那么第一个元素取的是src[0],而第二个元素则取的是src[4],第三个则是src[8],而写入到局部存储器中则是按次序前后相继写进去的。也就是第一个元素写到dst[0],第二个元素写到dst[1],第三个写到dst[2],对于从局部存储器异步拷贝到全局存储器的函数中,则增加了dst_stride参数,这个参数指明了从局部存储器所取出的元素写入到全局存储器中前一个与后一个相隔多少个元素。这意味着数据从局部存储器中取出时是逐个相继取出的,而写入到全局存储器时是以dst_stride跨度去写的。例如,取出的第一个数据是src[0],第二个数据是src[1],第三个数据是src[2],如果dst_stride是4,那么写入的第一个数据是dst[0],第二个数据是dst[4],第三个数据则是dst[8]。
下面我们根据上面异步拷贝操作的例子进行稍许修改来给出带有跨度的异步拷贝操作的使用实例。下面,我们将上述代码中最后的校验部分替换为如下代码:

//做数据校验
bool isOK = true;
for(int i = 0; i 〈 contentLength / sizeof(int); i++)
{
    //跨4个int元素进行比较校验
    if((i & 3) == 0 && pHostBuffer[i] * 2 != pDeviceBuffer[i])
}
    {
        isOK = false;
        break;
    }
puts(isOK ? "OK" : "NG");

由于我们这里写回到全局存储空间已经是间隔了4个int元素,所以我们在进行比较校验时也必须间隔4个int元素进行比较,否则结果肯定会失败。下面给出kernel程序代码:

__kernel void kernel_test(__global int *pDst,
{
                              __global int *pSrc1, __global int *pSrc2)
    local int tmpBuffer[GROUP_NUMBER_OF_WORKITEMS];
    const size_t address_offset = get_group_id(0) *
                                        GROUP_NUMBER_OF_WORKITEMS;
    //这里使用了异步拷贝操作,
    //将全局存储器的数据拷贝到相应的当前工作组的局部存储器中
    //这里,函数原型中的gentype就是int类型
    //一共拷贝GROUP_NUMBER_OF_WORKITEMS / 4个元素
    //从全局存储器获取的两个元素之间跨4个元素(即4个int大小)
    event_t event = async_work_group_strided_copy(tmpBuffer,
                                            &pSrc1[address_offset],
                                    GROUP_NUMBER_OF_WORKITEMS / 4,
                                                                4, 0);
    //这里抽空可以做其他事情,比如确定当前工作组中当前工作项的id等
const int index = get_local_id(0);
    //等待拷贝结束
    wait_group_events(1, &event);
    //我们将对应元素乘以2,再写回局部存储器
tmpBuffer[index] *= 2;
//我们确保所有工作项都在调用后续的异步操作之前能够都完成对数据的修改操作
    barrier(CLK_LOCAL_MEM_FENCE);
    //将当前局部存储器中的数据再拷贝回对应的全局存储器中,并且相邻两个元素之间跨4个元素
    event = async_work_group_strided_copy(&pDst[address_offset],
                                                  tmpBuffer,
                                                  GROUP_NUMBER_OF_WORKITEMS / 4,
                                                  4, 0);
    //等待拷贝结束
    wait_group_events(1, &event);
}

赞(0)
未经允许不得转载:极客笔记 » 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映射