时间戳

首页 > 解决方案 > 在 OpenCL 1.2 上实现堆栈推送的正确方法是什么？

algorithm - 在 OpenCL 1.2 上实现堆栈推送的正确方法是什么？

问题描述

问题

假设多个工作项要附加到全局堆栈：

void kernel(__global int* stack) {
    ... do stuff ...
    push(stack, value);
    ... do stuff ...
    return y;
}

理想的是，内核运行后，stack包含所有value推送给它的内容。顺序无所谓。在 OpenCL 1.2 中执行此操作的正确方法是什么？

我试过的

一个明显的想法是使用atomic_inc来获取长度并写入它：

void push(__global int* stack, int val) {
    int idx = atomic_inc(stack) + 1; // first element is the stack length
    stack[idx] = val;
}

但我推测在相同的内存位置上分别调用所有工作项atomic_inc会破坏并行性。一个单独的想法是只写入一个大于工作项数量的临时数组：

void push(__global int* stack, int val) {
    stack[get_global_id(0)] = val;
}

这会给我们留下一个稀疏的值数组：

[0, 0, 0, 7, 0, 0, 0, 2, 0, 0, 3, 0, 0, 0, 9, 0, 0, ...]

然后可以使用"stream compaction"对其进行压缩。因此，我想知道这些想法中哪一个是最有效的，以及是否可能存在我不知道的第三种选择。

标签： algorithmdata-structuresstackopencl

解决方案

我不能在这里给你一个明确的答案，但我可以提出一些尝试的建议——如果你有资源，尝试实现其中的一个以上，并在你所有不同类型的 OpenCL 实现中分析它们的性能'正计划部署。您可能会发现不同的解决方案在不同的硬件/软件上执行不同。

  1. 在本地内存中为每个工作组创建一个堆栈（显式地或在生成所有值后通过压缩），并且仅将全局堆栈增加每个工作组的计数，并将整个本地堆栈复制到全局堆栈中。这意味着每个工作组只有一个全局原子添加。当然，对于大型团体来说效果更好。
  2. 在幼稚方法中，您最大的原子争用来源将来自同一工作组中的项目。因此，您可以为每个工作组创建与项目一样多的堆栈，并让组中的每个项目提交到其“自己的”堆栈。在此之后，您仍然需要一个压缩步骤将其全部合并到一个列表中。如果您尝试这样做，请改变组大小。我不确定当前的 GPU 在多大程度上受到错误共享的影响（原子锁定整个缓存行，而不仅仅是那个词），因此您需要检查和/或尝试内存中堆栈计数器之间的不同间隙。
  3. 将所有结果写入固定偏移量（基于全局 id）一个足够大的数组以捕获最坏的情况，并将一个单独的压缩内核排队，该内核将结果后处理成一个连续的数组。
  4. 不要为结果的紧凑表示而烦恼。相反，使用稀疏数组作为下一阶段计算的输入。下一阶段的工作组可以将稀疏数组的固定子集压缩到本地内存中。完成后，每个工作项都在压缩数组的一项上工作。在内核内部迭代，直到全部处理完毕。这工作的好坏取决于数组中稀疏项的统计分布的可预测性，以及您选择的工作组大小以及每个工作组处理的稀疏数组的数量。这个版本还避免了到主机处理器的往返。
  5. 特别是在英特尔 IGP 上，我听说具有可变输出数量的 DirectX/OpenGL/Vulkan 几何着色器表现异常出色。如果您可以以几何着色器的格式编写算法，那么如果您以这些设备为目标，这可能值得一试。对于 nvidia/AMD，不要为此烦恼。

可能还有其他选择，但这些应该会给你一些想法。

推荐阅读