OpenCL 子内核存储器可见性

设备队列引入了一个非常重要的问题:父子内核存储器之间的一致性如何保证,即子内核访问父内核分配的存储器会产生怎样的行为。依据存储器的生命周期和是否可写,这个问题可以依据存储器类型的不同而分别描述。

全局存储器

父内核和子内核访问全局存储器有相干性,子内核与父内核之间有弱一致性。在执行过程中有两个时间点子内核和父内核的存储器是完全一致性的:一是父内核创建子内核时;二是父内核在子内核完成后调用同步API时。在启动入队子内核之前,父内核对全局存储器的操作对子内核是可见的,所有父内核对全局存储器的操作都会影响到子内核。父内核在子内核完成调用同步后,子内核对全局存储器的操作父内核可见。

常量存储器

常量存储器在内核执行期间不能被修改。所有的常量存储器必须在主机端第一次入队内核时进行初始化。常量对所有内核都是可见的,在父内核和子内核中常量都是不变的。

局部存储器

局部存储器是工作项私有的,当前工作组以外的工作项都不能访问。把指向工作组内的局部存储器的指针传递给子内核,结果都是未定义的。例子如下:

kernel void foo(global int *a, local int *lptr, ...)
{
    ……
    enqueue_kernel(get_default_queue(),
                      CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
                      ndrange,
                      ^
    {
        size_t id = get_global_id(0);
        local int *p = lptr; //未定义行为
    };
}

把指向父内核中局部存储器作为参数传入子内核,程序编译不会报错,但是执行结果是未知的。
但是在父内核中,我们可以为子内核分配所需要的局部存储器。如下例子展示了如何在父内核中给子内核分配一个或多个全局存储器大小:

kernel void my_func_A_local_arg1(
    global int *a, local int *lptr, ...)
{
    ……
}
kernel void my_func_A_local_arg2(
    global int *a, local int *lptr1, local float4 *lptr2, ...)
{
    ...
}
kernel void my_func_B(global int *a, ...)
{
    ...
    ndrange_t ndrange = ndrange_1d(...);
    uint local_mem_size = compute_local_mem_size();
    enqueue_kernel(get_default_queue(),
                      CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
ndrange,
                      ^(local int *p)
    {
        my_func_A_local_arg1(a, p, ...);
    },
local_mem_size);
}
kernel void my_func_C(global int *a, ...)
{
    ...
    ndrange_t ndrange = ndrange_1d(...);
    void (^my_blk_A)(local int *, local float4 *) =
        ^(local int *lptr1, local float4 * lptr2)
    {
        my_func_A_local_arg2(a, lptr1, lptr2, ...);
    };
    //calculate local memory size for lptr
    //argument in local address space for my_blk_A
    uint local_mem_size = compute_local_mem_size();
    enqueue_kernel(get_default_queue(),
                      CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
                      ndrange,
                      my_blk_A,
    local_mem_size, local_mem_size * 4);
}

代码中粗体字部分标出在父内核中如何分配子内核局部存储器大小。子内核函数局部存储器参数作为块的输入参数,分配的局部存储器具体大小由enqueue_kernel()函数最后一个参数确定。

私有存储器

私有存储器是工作项的私有存储空间,当前工作项之外的其他工作项都对它不可见。把用指向私有存储器变量的指针作为子内核传入参数是非法的。

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