【问题标题】:Can be__shfl_xor replaced in order to run on sm_21?可以替换__shfl_xor 以便在sm_21 上运行吗?
【发布时间】:2017-09-01 13:12:43
【问题描述】:

尝试运行这个: https://github.com/Celebrandil/CudaSiftNVS4200M 上,即 sm_21,而不是所需的 sm_35。 运行上述项目的唯一问题是这段代码(cudaSiftD.cu:205):

for (int i=1;i

是否有可能的等效代码?

【问题讨论】:

  • 是的,如果你愿意写的话。
  • 几乎所有你可以用shuffle操作做的事情都可以用共享内存操作来完成,这也允许线程间通信。我并不是说实现是相同的,只是有一个使用共享内存的“可能的等效代码”。
  • @talonmies 此评论对 OP 有何帮助?这是一个重要的问题,因为我不认为 shuffle 内在函数是 cuda 的一个简单特性。

标签: cuda gpgpu sift


【解决方案1】:

嗯,几乎所有 CUDA 内在函数都可以替换,所以我将您的问题解释为

__shfl_xor 可以在 SM_21 GPU 上便宜替换吗?

答案是:并非如此;你会受到惩罚。正如@RobertCrovella 的评论所暗示的,你最好的选择是使用共享内存:

  • 每个通道将其数据写入共享内存中的一个位置(使这些连续的 4 字节大小的值避免bank conflicts
  • 执行某种同步(可能您必须__syncthreads()
  • 每个通道都从共享内存位置读取它想要写入其值的通道。

我没有拼出代码来不为你带走乐趣:-)

编辑: 虽然 shuffle 的执行更复杂,但至少在语义上它仍然是对寄存器的操作;它不需要同步。所以共享内存的替代方案会更慢。

【讨论】:

  • 我不会认为 shuffle 是一个时钟周期,原因有两个:1)在多处理器上每个周期有 32 个可发布的 shuffle [docs.nvidia.com/cuda/cuda-c-programming-guide/…,2)shuffle 操作由缓存执行管理共享内存。从本质上讲,使用 shuffle 的性能大约是共享内存的两倍 - 请参阅 [on-demand.gputechconf.com/gtc/2013/presentations/…
  • @FlorentDUGUET:已编辑以反映您的评论。您的链接不起作用,我认为您的括号中有一些错字。
  • @FlorentDUGUET:评论 +1 指出 shuffle 执行涉及访问共享内存的机制。
  • 是的,我知道有一个“价格”(也就是惩罚)。问题是 - 我不熟悉 CUDA API,所以我自己编写代码不是我的选择,至少现在是这样。我只是在寻找可以与 Python 一起使用的 SIFT 的 CUDA 加速实现。我找到了,但遇到了问题中描述的问题。无论如何,谢谢你试图帮助我!
【解决方案2】:

如果问题更多是关于如何用与 sm_21 兼容的代码替换此 sn-p 代码,您可能需要关注 CUB,即块缩减部分 here。模板参数之一是您设备的架构。

__CUDA_ARCH__ 宏可以帮助您选择最合适的实现,请参阅here

【讨论】:

    猜你喜欢
    • 2012-04-19
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 2010-09-19
    • 2021-05-04
    • 2012-06-20
    • 2012-05-10
    • 1970-01-01
    相关资源
    最近更新 更多