__shfl_down与__shfl_down_sync执行结果不一致技术咨询

哇，这个问题的根源其实很简单——你把__shfl_down_sync的参数顺序搞反啦！

先明确两个函数的核心参数差异，这是问题的关键：

• __shfl_down的原型是：

  template <typename T>
  __device__ T __shfl_down(T var, unsigned int delta, int width=warpSize);
  

  参数顺序是：要交换的变量var → 偏移量delta → 逻辑warp大小width。
• __shfl_down_sync的原型是：

  template <typename T>
  __device__ T __shfl_down_sync(unsigned int mask, T var, unsigned int delta, int width=warpSize);
  

  它比前者多了第一个线程掩码参数mask，后面才是var、delta、width。

你的第一个内核调用__shfl_down(temp, offset,32)是完全正确的：对应var=temp、delta=offset、width=32，实现了warp内的归约逻辑。你的block是64线程，分成两个独立warp（0-31和32-63），每个warp有32个线程，初始值都是2.0，归约后总和为2*32=64.0，所以所有线程最终输出64.0。

但第二个内核里，你写成了__shfl_down_sync(temp, offset,32)，这就完全偏离了预期：你把temp当成了掩码参数，把offset当成了要交换的变量，把32当成了偏移量delta！

我们拆解这个错误调用的行为：

1. 每次循环中，实际执行的是temp += __shfl_down_sync(temp, offset, 32)，这里的var是整数offset（会被隐式转为double），delta=32，width=32。
2. 根据CUDA的shfl规则，当lane + delta >= width时，会返回当前线程自己的var值。因为delta=32、width=32，所有线程的lane号加32都大于等于32，所以每次取的都是当前线程的var——也就是offset的数值。
3. 你的循环offset依次是16、8、4、2、1，初始temp=2.0，累加后就是2 + 16 +8 +4 +2 +1 =33，这正好是你看到的第二个内核的输出结果。

要修复这个问题，只需要把__shfl_down_sync的参数顺序改对，传入正确的掩码（比如用0xffffffff表示当前warp所有线程，或者用__activemask()获取活跃线程掩码）：

temp += __shfl_down_sync(0xffffffff, temp, offset, 32);

修改后，第二个内核的输出就会和第一个一致，所有线程都输出64.0。

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