深入了解__shfl__sync()中的第一个参数掩码

Question

以下是广播变量的测试代码：

#include <stdio.h>
#include <cuda_runtime.h>
__global__ void broadcast(){
int lane_id = threadIdx.x & 0x1f;
int value = 31 - lane_id;
//let all lanes within the warp be broadcasted the value 
//whose laneID is 2 less than that of current lane
int broadcasted_value = __shfl_up_sync(0xffffffff, value, 2)
value = n;
printf("thread %d final value = %dn", threadIdx.x, value);
}
int main() {
broadcast<<<1,32>>>();
cudaDeviceSynchronize();
return 0;
}

实际上，这个问题与本页的问题相同。无论我修改了什么掩码(例如0x00000000、0x00000001等(，混洗的结果都没有变化。那么，如何正确理解掩码的效果呢？

Answer 1

掩码参数在执行请求的混洗操作之前，对于用1位标识的曲速通道，强制曲速重新聚合(假设这种重新聚合是可能的，即条件编码不阻止。如果条件编码阻止，则代码是非法的，并探索未定义的行为-UB(。

对于已经聚合并处于活动状态的经线，它没有任何效果如果掩码参数设置为零，则不会阻止通道参与混洗操作它也不会强制非活动车道参与(非活动车道将是被条件编码排除在外的车道(。

由于您的代码没有条件行为，因此没有理由相信会缺乏收敛性，因此无论您的掩码参数如何，行为都不会发生变化。

这并不意味着将掩码指定为0是正确的。如果您希望泳道参与，但没有在掩码中将其相应的位设置为1，则您的代码是非法的，并且在发生扭曲发散的情况下，您可能正在探索UB。

关于面具的其他描述，这里已经有很多答案了。

1。2.3.4.5.

你可能提出的任何后续问题都有可能在其中一个问题中得到回答。