shfl.sync.idx 什么时候快？

Question

使用

.idx

 的 

shfl.sync 选项，可以在单个 warp 中的线程之间任意排列寄存器。希望通过使用

shfl.sync

，您可以避免从共享内存中存储然后加载数据。

但是，根据this HN comment，听起来

shfl.sync

有时会经过共享内存（或其他一些慢速路径），具体取决于具体的shuffle配置：

在 CUDA PTX 中，有一条名为“shfl.sync.idx”的指令，它允许您在线程之间任意调整数据。然而，只有遵循各种规则，这种洗牌才会有效：避免银行冲突等等。

我在网上查了一下，但找不到任何资源来说明到底哪些洗牌速度快。

将

shfl.sync

与

.idx

一起使用时，总是很快吗？它总是通过寄存器吗？如果没有，如果我想要

shfl.sync

和

.idx

走上快速道路，我需要遵循什么规则？

相关：https://stackoverflow.com/a/76123988/3154996

Answer 1

shfl

指令在扭曲中的不同线程之间执行收集操作。

其工作方式是 (A) 每个线程公开其自己的 32 位值，并且 (B) 从 A 中公开的菜单中选择 32 个值之一。
这总是通过寄存器起作用，共享内存不被触及。

缺点是只能在

idx

中指定来源，而不能指定目的地。这意味着您不能轻松地进行分散操作，只能进行聚集操作。

idx

变体是最通用的，每个线程都可以指定不同的索引。

您可以使用以下代码轻松计算变体的时间：

__global__ void time_shfl(int a) { //parameter to stop the optimizer from eliminating code
  const auto a = i * threadIdx.x;
  const auto source = 31 - threadIdx.x; //reverse the order
  const auto StartTime = clock64();
  const auto b = shfl_sync(-1u, a, source, 32);
  const auto EndTime = clock64();
  printf("T:%i, a = %i, b = %i, time = %i\n", threadIdx.x, a, b, int(EndTime - StartTime);
}

int main() {
  time_shfl<<<1, 32>>>(10);
  cudaDeviceSynchronize();
}

WIP，后续时间安排。