OpenCL中工作项和全局内存之间的内存传输？

Posted 2023-02-16

【中文标题】OpenCL中工作项和全局内存之间的内存传输？

【英文标题】：Memory transfer between work items and global memory in OpenCL?

【发布时间】：2014-04-10 15:30:18

【问题描述】：

我对工作项和全局内存之间的数据传输方式有一些疑问。让我们考虑以下效率极低的内存绑定内核。

__kernel void reduceURatios(__global myreal *coef, __global myreal *row, myreal ratio)

    size_t gid = get_global_id(0);//line no 1

    myreal pCoef = coef[gid];//line no 2
    myreal pRow = row[gid];//line no 3

    pCoef = pCoef - (pRow * ratio);//line no 4
    coef[gid] = pCoef;//line no 5

工作组中的所有工作项是否开始执行第 1 行 同一时间？ 工作组中的所有工作项是否开始执行第 2 行 同一时间？ 假设工作组中的不同工作项完成执行行 没有 4 在不同的时间。做早完成的等待，以便， 所有工作项同时将数据传输到全局内存 在第 5 行？ 所有工作项是否同时退出计算单元，以便 提前完成的工作项必须等到所有工作项都完成 执行完毕？ 假设每个内核必须从全局内存执行 2 次读取。是吗 最好一个接一个地执行这些语句，或者是 最好在两次读取之间执行一些计算语句 处决？ 上面显示的内核是用于 GPU 的内存。有没有办法通过 哪些性能可以提高？ 是否有避免内存限制的一般准则？

【参考方案1】：

在下面找到我的答案：（感谢Sharpneli 对AMD GPU 和warp 的好评）

通常是。但取决于硬件。您不能直接期望这种行为并在这种“有序执行”上设计您的算法。这就是存在barriers 和mem_fences 的原因。例如，某些 GPU 仅按顺序执行 WG 的 WI 的子集。在 CPU 中，它们甚至有可能完全无序地运行。 同答案 1。 如答案 1 所示，它们不太可能在不同的时间完成，所以是的。但是您必须记住，这是一个很好的功能，因为 1 次大写内存比大量小写更有效。 通常是的（也请参见答案 1） 最好将读取插入操作，但编译器已经考虑到这一点并重新排序操作顺序以隐藏读取/写入效果的延迟。当然，编译器永远不会移动可以改变结果值的代码。除非您手动禁用编译器优化，否则这是 OpenCL 编译器的典型行为。 不，从内核的角度来看，它无法以任何方式进行改进。 一般规则是，输入的每个存储单元都被多个 WI 使用？ 否（1 个全局->1 个私有） （这是您的内核在问题中的情况） 那么那个内存是全局->私有的，没办法改进，不要用本地内存，浪费时间。 是（1 个全局 -> X 个私有） 尝试先将全局内存移至本地内存，然后将每个 WI 直接从本地读取到私有。根据重用量（可能只有 2 个 WI 使用相同的全局数据），如果计算量已经很高，甚至可能不值得。您必须考虑额外内存使用和全局访问增益之间的权衡。对于图像处理来说，这通常是一个好主意，但对于其他类型的处理则不是这样。

注意：如果您尝试写入全局内存，同样的过程也适用。在写入全局之前，由许多 WI 在本地内存中操作总是更好。但是如果每个WI写入全局唯一的地址，那么直接写入。

【讨论】：

实际上 1 和 2 的答案在很大程度上取决于。例如，256 个 AMD 硬件工作组有 64 个项目执行一行，甚至可能更多，然后该工作组中的下一个 64 才开始执行。必须在工作组中使用本地内存栅栏是有充分理由的。

是的，您是对的，您实际上无法预测所有线程将一起运行。例如，因为它们可以进一步划分为小的并行工作（即：64 WI 一起分成 4 个块以形成 256 WG）。这就是栅栏和障碍存在的原因。我会更新我的答案。