shfl.sync.idx何时能实现高效运行？

带.idx的shfl.sync并非始终高效，也不会始终通过寄存器直接传输——当洗牌配置不符合硬件要求时，确实会 fallback 到共享内存这类慢路径。以下是需要遵循的关键规则，以及相关细节：

• 避免存储体冲突：GPU的warp洗牌硬件按存储体（bank）划分，每个线程对应固定存储体。如果洗牌操作中多个线程同时访问同一个存储体，就会引发冲突，导致硬件不得不通过共享内存完成数据传输。

  注：PTX层面的shfl.sync.idx不直接暴露存储体编号，但可通过线程ID的映射逻辑规避——比如确保目标线程ID的低N位（N取决于GPU架构的存储体数量，如Volta及以后是32个存储体，对应低5位）不重复访问同一值。

• 遵循架构特定的对齐要求：不同GPU架构对洗牌操作的对齐规则有差异：
  • Kepler及更早架构：仅支持有限的洗牌模式（如上下移动、广播），shfl.sync.idx这类任意洗牌会直接走共享内存；
  • Maxwell到Turing架构：支持任意洗牌的快速路径，但要求目标线程ID的计算是线程ID的线性函数（比如dst_tid = (tid + offset) % warp_size这类简单映射）；
  • Ampere及以后架构：对任意洗牌的支持更灵活，但仍需避免存储体冲突才能维持寄存器级传输的效率。
• 禁止跨warp访问：shfl.sync.idx仅能在单个warp内操作，若指定的目标线程ID超出当前warp范围（0~31），会触发未定义行为，同时大概率走慢路径。

补充细节 #

对于Ampere及更新的架构，shfl.sync.idx的快速路径支持绝大多数任意洗牌场景，只要不触发存储体冲突。而在旧架构（如Maxwell）中，只有当洗牌模式是**“线程间的偏移固定且连续”**时，才会走寄存器级的快速路径——比如每个线程从tid + k（k为固定偏移）的线程获取数据，这类模式可被硬件直接优化。如果是完全随机的线程间置换（比如线程0从线程17取，线程1从线程9取等无规律映射），旧架构就会不得不使用共享内存完成数据交换。

内容的提问来源于stack exchange，提问作者Elliot Gorokhovsky