深入了解__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 = %d\n", threadIdx.x, value);
}

int main() {
    broadcast<<<1,32>>>();
    cudaDeviceSynchronize();
    return 0;
}

实际上，此问题与this page相同。无论我修改了遮罩（例如0x00000000，0x00000001等），改组的结果都没有改变。那么，如何正确理解口罩的作用呢？

这里是广播变量的测试代码：#include

#include __global__ void broadcast（）{int lane_id = threadIdx.x＆0x1f; int值= 31-...

Answer 1

掩码参数在执行请求的混洗操作之前，对于以1位标识的翘曲通道强制进行翘曲收敛（假设这种收敛是可能的，即不受条件编码的限制。如果被条件编码阻止，则您的代码是非法的，并探索未定义的行为-UB）。

对于已经收敛并处于活动状态的经线，它没有任何作用。

如果mask参数设置为零，它不会阻止泳道参与随机播放操作。