【问题标题】: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 谢谢这很有用

标签: opencl bank-conflict


【解决方案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
  • 酷。谢谢,蝉。
猜你喜欢
  • 2015-04-07
  • 2012-10-02
  • 2011-01-18
  • 2015-07-20
  • 2013-02-07
  • 1970-01-01
  • 1970-01-01
  • 2011-03-31
  • 1970-01-01
相关资源
最近更新 更多