varp primitive函数中的mask问题

  • 3 replies
  • 114 views
varp primitive函数中的mask问题
« 于: 十一月 10, 2019, 02:17:40 am »
我写了如下测试代码:

#include <stdio.h>

const unsigned FULL_MASK = 0xffffffff;

void __global__ test_warp_primitives(void)
{
    int tid = threadIdx.x;

    unsigned mask = __ballot_sync(FULL_MASK, tid >= 1);
    if (tid == 0) printf("FULL_MASK     = %x\n", FULL_MASK);
    if (tid == 0) printf("mask          = %x\n", mask);

    int v1 = tid;
    v1 = __all_sync(FULL_MASK, tid);
    if (tid == 0) printf("all_sync:     %d\n", v1);

    int v2 = tid;
    v2 = __all_sync(mask, tid);
    if (tid == 0) printf("all_sync:     %d\n", v2);

    int v3 = tid;
    v3 = __shfl_sync(FULL_MASK, tid, 2);
    if (tid == 0) printf("shfl (FULL_MASK): ");
    printf("%d ", v3);
    if (tid == 0) printf("\n");

    int v4 = tid;
    v4 = __shfl_sync(mask, tid, 2);
    if (tid == 0) printf("shfl (mask):      ");
    printf("%d ", v4);
    if (tid == 0) printf("\n");
}

int main(int argc, char **argv)
{
    test_warp_primitives<<<1, 32>>>();
    cudaDeviceSynchronize();
    return 0;
}

输出为:

FULL_MASK     = ffffffff
mask          = fffffffe
all_sync:     0
all_sync:     1
shfl (FULL_MASK): 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2
shfl (mask):      2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2

我不理解最后一行的输出。最后一行的输出是如下变量的结果:
int v4 = __shfl_sync(mask, tid, 2);
这里,我给__shfl_sync函数传入了一个掩码mask=fffffffe。按道理,这将会是第0号线程不参与计算。所以,我以为最后一行的输出结果应该是:
shfl (mask):      0 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2

另外,程序的如下两行输出是我能理解的:
all_sync:     0
all_sync:     1
第一行用是用FULL_MASK的结果:有一个线程的值是0,所以__all_sync给出0;
第二行是用mask的结果:所有“参与”线程的值是非零,所以__all_sync给出1。

所以,我对程序最后一行的输出感到很困惑,还请大家解惑。谢谢!

Re: varp primitive函数中的mask问题
« 回复 #1 于: 十一月 10, 2019, 02:23:34 am »
补充一下我的测试环境:

CUDA 9.0
CC=5.0的Quadro K2200
编译命令:nvcc -arch=sm_50 xxx.cu

Re: varp primitive函数中的mask问题
« 回复 #2 于: 十一月 11, 2019, 12:37:12 pm »
bruce你好,

对于不在mask中的线程,结果将是未定义的。此时出现任何结果都有可能。请不要使用未定义的行为(即,你的0号线程不参与shfl,却要求使用shfl的结果),这样是不安全的。

Re: varp primitive函数中的mask问题
« 回复 #3 于: 十一月 11, 2019, 05:04:48 pm »
Thank you very much! It's a simple and clear answer!