如何在 OpenCL 中基于数据同步（特定）工作项？

Question

上下文：

需要模拟相关分立元件（复杂电子电路）的网络。因此每个组件接收来自其他几个组件的输入并输出到其他几个组件。

预期的设计是 kernel，带有定义它应代表哪个组件的配置参数。电路的每个组件都由一个工作项表示，并且所有电路都将适合一个工作组（或者将对电路进行适当的拆分，以便每个工作组可以将所有组件作为工作项进行管理） .

问题：

是否可能，万一如何呢？让一些工作项等待其他工作项数据？工作项生成数组的输出（在数据驱动的位置）。另一个工作项需要等待这种情况发生才能开始处理。网络没有循环，因此，单个工作项不可能运行两次。

尝试次数：

在下面的示例中，每个组件最多可以有一个输入（为了简化），使电路成为一棵树，其中电路的输入是根，3 个输出是叶子。

inputIndex 通过为每个组件指示其他组件为其提供输入来对这棵树进行建模。第一个组件将自身作为输入，但内核管理这种情况（为了简化）。

result保存每个分量的结果（电压、强度等）

inputModified 表示给定组件是否已经计算出他的输出。

// where the data come from (index in result)
constant int inputIndex[5]={0,0, 0, 2, 2};

kernel void update_component(
    local int *result, // each work-item result. 
    local int *inputModified // If all inputs are ready (one only for this example)
) {

    int id = get_local_id(0);
    int size = get_local_size(0);
    int barrierCount = 0;

    // inputModified is a boolean indicating if the input is ready
    inputModified[id]=(id!=0 ? 0 : 1);

    // make sure all input are false by default (except the first input).
    barrier(CLK_LOCAL_MEM_FENCE); 


    // Wait until all inputs are ready (only one in this example)
    while( !inputModified[inputIndex[id]] && size > barrierCount++)
    {
        // If the input is not ready, wait for it
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    // all inputs are ready, compute output
    if (id!=0) result[id] = result[inputIndex[id]]+1;
    else result[0]=42;

    // make sure any other work-item depending on this is unblocked
    inputModified[id]=1;

    // Even if finished, we needs to "barrier" for other working items.
    while (size > barrierCount++)
    {
        barrier(CLK_LOCAL_MEM_FENCE);
    }
}

此示例对 N 个组件有 N 个障碍，使其比顺序解决方案更糟糕。

注意：这只是内核，因为 C++ 最小主机相当长。有需要的话，我可以想办法加上。

问题：

是否有可能通过内核本身有效地让不同的工作项等待其他工作项提供它们的数据？或者什么解决方案是有效的？

这个问题（对我来说）解释起来并不简单，而且我远不是 OpenCL 专家。请耐心等待，如果有任何不清楚的地方，请随时询问。

Answer 1

来自屏障的文档

https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/barrier.html

If barrier is inside a loop, all work-items must execute the barrier for each iteration of the loop before any are allowed to continue execution beyond the barrier.

但是内核中的一个while循环（包含一个barrier）有这样的条件：

inputModified[inputIndex[id]]

这可能会改变其线程 ID 的行为并导致未定义的行为。再说之前还有一个关卡

barrier(CLK_LOCAL_MEM_FENCE);

已经同步了工作组中的所有工作项，因此 while 循环是多余的，即使它有效。

最后一个屏障循环也是多余的

while (size > barrierCount++)
{
    barrier(CLK_LOCAL_MEM_FENCE);
}

当内核结束时，它会同步所有工作项。

如果您打算向工作组外的工作项发送一些消息，那么您只能使用原子变量。即使在使用原子时，您也不应假设任何两个工作项之间存在任何 working/issuing 顺序。

你的问题

how? to have some work-items waiting for other work-items data? A work-item generate an output to an array (at a data-driven position). Another work-item needs to wait for this to happens before to start making it processing. The net has no loops, thus, it is not possible that a single work-item needs to run twice.

可以用 OpenCL 2.x 特性 "dynamic parallelism" 来回答，它可以让工作项在内核中生成新的 workgroups/kernels。它比等待自旋等待循环更有效，并且绝对比依赖 GPU 支持的运行中线程数更独立于硬件（当 GPU 无法处理那么多运行中线程时，任何自旋等待都会死锁，线程顺序无关紧要）。

当你使用屏障时，你不需要通知其他线程"inputModified"。结果数据在工作组中已经可见。

如果你不能使用 OpenCL v2.x，那么你应该使用 BFS 处理一棵树：

为顶级节点启动 1 个工作项
对其进行处理并准备 K 个输出并将它们推入队列
结束内核
启动 K 个工作项（每个从队列中弹出元素）
处理它们并准备 N 个输出并将它们推入队列
结束内核
重复直到队列没有更多元素

内核调用次数等于树的最大深度，而不是节点数。

如果您需要比 "kernel launches" 更快的同步，则对整个树使用单个工作组，使用屏障而不是内核召回。或者，在 CPU 上处理前几个步骤，有多个子树并将它们发送到不同的 OpenCL 工作组。也许在 CPU 上计算直到有 N 个子树，其中 N=GPU 的计算单元对于基于工作组屏障的子树更快的异步计算可能更好。

还有一种无障碍、无原子、单内核调用的方式。从底部开始树并向上移动。

将所有最深层次的子节点映射到工作项。将它们中的每一个移动到顶部，同时在它们的私有内存/其他一些快速内存中记录它们的路径（节点 ID 等）。然后让他们通过记录的路径自上而下地遍历，在移动中计算，没有任何同步，甚至没有原子。这比 barrier/kernel-call 版本的工作效率低，但由于缺少屏障并且处于完全异步路径上，因此速度应该足够快。

如果树的深度为 10，这意味着要保存 10 个节点指针，而不是私有寄存器。如果树深度约为 30 40，则在每个工作组中使用较少线程的本地内存；如果更多，则分配全局内存。

但是您可能需要根据工作项的空间性/树的拓扑结构对工作项进行排序，以使它们在减少分支的情况下更快地协同工作。

这种方式我觉得最简单，建议你先试试这个无障碍版本

如果您只需要每个工作项的数据可见性而不是组或内核，请使用 fence：https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/mem_fence.html

如何在 OpenCL 中基于数据同步（特定）工作项？

How to synchronize (specific) work-items based on data, in OpenCL?

gpgpu

opencl

上下文：

问题：

尝试次数：

问题：