二维内核中的银行冲突

Posted 2023-04-15

【中文标题】二维内核中的银行冲突

【英文标题】：Bank conflict in 2D kernel

【发布时间】：2014-12-18 06:31:08

【问题描述】：

假设我们的硬件有 32 个 4 字节宽的存储体。我们有一个大小为 32 的一维内核和一个本地一维整数数组。

然后，确保每个连续线程访问数组中连续的内存位置应该避免银行冲突。

但是，假设我们有一个 8 x 4 2D 内核和相同的 1D 数组。如何确保没有银行冲突？我们如何定义二维数组的“连续线程”？

【问题讨论】：

简单地将线程 ID 从 2D 线性化到 1D。

谢谢，蝉。你的意思是 get_local_id(0) + get_local_id(1) * 4 吗？

This post 及其答案有助于了解如何定义 2D 工作组中的连续线程。

@Farzad 谢谢这很有用

【参考方案1】：

您可以获得与在 1D 情况下使用 get_global_id(0) 在 2D 情况下使用此代码获得的相同的全局工作项 ID：

get_global_id(1) * get_global_size(0) + get_global_id(0);

如果您想在工作组中获取本地工作项 ID，只需将全局变量更改为本地变量即可。

【讨论】：

谢谢，bl0z0。这在任何地方都有记录吗？

我自己弄明白了，因为我处理过很多存储为一维数组的图像。但是您也可以从 Fazard 给您的链接中找到它。

IIRC 它位于§3.2 Execution Model

酷。谢谢，蝉。