深入了解 __shfl__sync() 中的第一个参数掩码
Insight into the first argument mask in __shfl__sync()
这里是广播变量的测试代码：
#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 = %d\n", threadIdx.x, value);
}
int main() {
broadcast<<<1,32>>>();
cudaDeviceSynchronize();
return 0;
}
实际上，这个问题和this page的问题是一样的。无论我修改掩码（例如 0x00000000、0x00000001 等），洗牌的结果都没有变化。那么，如何正确认识面膜的功效呢？
掩码参数强制扭曲重新收敛，对于用 1 位标识的扭曲通道，在执行请求的洗牌操作之前（假设这种重新收敛是可能的，即不被条件编码阻止。如果被条件编码阻止，您的代码是非法的，正在探索未定义的行为 - UB）。
对于已经收敛并激活的曲速通道，它没有效果。 如果掩码参数设置为零，它不会阻止通道参与洗牌操作。它也不会强制非活动通道参与（非活动通道将是被条件排除的通道编码）。
由于您的代码没有条件行为，因此没有理由相信会缺乏收敛，因此无论掩码参数如何，行为都不会发生变化。
这并不意味着将掩码指定为 0 是正确的。如果您希望通道参与但未在掩码中将其相应位设置为 1，则您的代码是非法的，并且您可能会在中探索 UB经线发散事件。
关于面膜的其他描述，这里已经有很多答案了。
5.
您可能有的任何后续问题可能已经在其中一个中得到了解答。
这里是广播变量的测试代码：
#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 = %d\n", threadIdx.x, value);
}
int main() {
broadcast<<<1,32>>>();
cudaDeviceSynchronize();
return 0;
}
实际上，这个问题和this page的问题是一样的。无论我修改掩码（例如 0x00000000、0x00000001 等），洗牌的结果都没有变化。那么，如何正确认识面膜的功效呢？
掩码参数强制扭曲重新收敛，对于用 1 位标识的扭曲通道，在执行请求的洗牌操作之前（假设这种重新收敛是可能的，即不被条件编码阻止。如果被条件编码阻止，您的代码是非法的，正在探索未定义的行为 - UB）。
对于已经收敛并激活的曲速通道，它没有效果。 如果掩码参数设置为零，它不会阻止通道参与洗牌操作。它也不会强制非活动通道参与（非活动通道将是被条件排除的通道编码）。
由于您的代码没有条件行为，因此没有理由相信会缺乏收敛，因此无论掩码参数如何，行为都不会发生变化。
这并不意味着将掩码指定为 0 是正确的。如果您希望通道参与但未在掩码中将其相应位设置为 1，则您的代码是非法的，并且您可能会在中探索 UB经线发散事件。
关于面膜的其他描述，这里已经有很多答案了。
您可能有的任何后续问题可能已经在其中一个中得到了解答。