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（）函数最后一个参数确定。