请教一下关于warp shuffle f4unction中的mask

[BigFaceBoy]

显卡：nvidia l40
测试代码：

#define SIZE 1024

__global__ void test_shfl_sync(int* x, int N){
    int tid = threadIdx.x;
    int val = tid ;

    x[tid] = __shfl_sync(0xffffffff, val, 0);
    __syncwarp();
}
int main(){
    int* x = (int *)malloc(sizeof(int) * SIZE);
    int* y = (int *)malloc(sizeof(int) * SIZE);

    for(int i = 0; i < SIZE; i++){
        x[i] = i;
        y[i] = 0;
    }

    int *d_x1;
    cudaMalloc(&d_x1, sizeof(int) * SIZE);
    cudaMemcpy(d_x1, x, sizeof(int) * SIZE, cudaMemcpyHostToDevice);

    int THREAD_NUM_PER_BLOCK = 64;
    int block_num = (SIZE + THREAD_NUM_PER_BLOCK - 1) / THREAD_NUM_PER_BLOCK;

    test_shfl_sync<<<block_num, THREAD_NUM_PER_BLOCK>>>(d_x1, SIZE); 

    cudaDeviceSynchronize();
    cudaMemcpy(y, d_x1, sizeof(int) * SIZE, cudaMemcpyDeviceToHost);

    for(int i = 0; i < 64; i++){
        printf("%d ", y[i]);
    }

    printf("\n");
}

打印是：0 0 ... 0 32 32 ... 32， 这个如预期，然后我修改mask为 0x0000ffff,

__global__ void test_shfl_sync(int* x, int N){
    int tid = threadIdx.x;
    int val = tid ;

    x[tid] = __shfl_sync(0x0000ffff, val, 0);
    __syncwarp();
}

打印还是：0 0 ... 0 32 32 ... 32， 这就搞不懂了

[BigFaceBoy]

我看两个kernel的汇编代码是一样的, 所以mask没有效果，但这是为什么呢

test_shfl_sync(int*, int):
   MOV R1, c[0x0][0x28] 
   S2R R0, SR_TID.X 
   MOV R3, 0x4 
   SHFL.IDX PT, R5, R0, RZ, 0x1f 
   IMAD.WIDE R2, R0, R3, c[0x0][0x160] 
   STG.E.SYS [R2], R5 
   NOP 
   EXIT