CUDA 9中shfl与shfl_sync的区别

Question

自CUDA 9以来，shfl指令已被弃用，应该用shfl_sync替换。但是它们的行为不同，我该如何替换它们呢？

代码示例：

__global__
static void shflTest(){
    int tid = threadIdx.x;
    float value = tid + 0.1f;
    int* ivalue = reinterpret_cast<int*>(&value);

    //use the integer shfl
    int ix = __shfl(ivalue[0],5,32);
    int iy = __shfl_sync(ivalue[0],5,32);

    float x = reinterpret_cast<float*>(&ix)[0];
    float y = reinterpret_cast<float*>(&iy)[0];

    if(tid == 0){
        printf("shfl tmp %d %d\n",ix,iy);
        printf("shfl final %f %f\n",x,y);
    }
}

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

输出：

shfl tmp 1084437299 5
shfl final 5.100000 0.000000

- dari

Answer 1

如果您阅读了CUDA 9 RC编程指南（第B.15节），并与您的CUDA 9 RC副本一起安装，您将看到新的__shfl_sync()函数有一个额外的mask参数，您没有考虑到它：

CUDA 8：

int __shfl(int var, int srcLane, int width=warpSize);

CUDA 9:

T __shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize);
              ^^^^^^^^^^^^^

这个掩码参数的期望也被指明了：

新的*_sync shfl内置函数需要一个掩码，表示参与调用的线程。每个参与线程所在的lane id表示为一个比特位，必须被设置，以确保它们在硬件执行内置函数之前被正确汇合。掩码中命名的所有未退出的线程都必须使用相同的掩码执行相同的内置函数，否则结果将是未定义的。

因此，如果我们修改你的代码来符合这个要求，就会得到预期的结果：

$ cat t419.cu
#include <stdio.h>

__global__
static void shflTest(int lid){
    int tid = threadIdx.x;
    float value = tid + 0.1f;
    int* ivalue = reinterpret_cast<int*>(&value);

    //use the integer shfl
    int ix = __shfl(ivalue[0],5,32);
    int iy = __shfl_sync(0xFFFFFFFF, ivalue[0],5,32);

    float x = reinterpret_cast<float*>(&ix)[0];
    float y = reinterpret_cast<float*>(&iy)[0];

    if(tid == lid){
        printf("shfl tmp %d %d\n",ix,iy);
        printf("shfl final %f %f\n",x,y);
    }
}

int main()
{
    shflTest<<<1,32>>>(0);
    cudaDeviceSynchronize();
    return 0;
}
$ nvcc -arch=sm_61 -o t419 t419.cu
t419.cu(10): warning: function "__shfl(int, int, int)"
/usr/local/cuda/bin/..//include/sm_30_intrinsics.hpp(152): here was declared deprecated ("__shfl() is deprecated in favor of __shfl_sync() and may be removed in a future release (Use -Wno-deprecated-declarations to suppress this warning).")

$ cuda-memcheck ./t419
========= CUDA-MEMCHECK
shfl tmp 1084437299 1084437299
shfl final 5.100000 5.100000
========= ERROR SUMMARY: 0 errors
$

- Robert Crovella

---

即使宽度不是32，用shfl_xx_sync(0xFFFFFFFF,...)替换所有的shfl_xx(...)是否安全？ - dari

是的，我理解。你知道 width 参数是如何工作的吗？ - Robert Crovella

是的，width参数将完整的warp宽度（32个lane）分成一组相等大小的2的幂子组，例如8个4组或4个8组。在每个子组内，相同的洗牌操作会发生，模除组宽度。因此，如果您有8个4组，但是您希望在每个8组中执行洗牌操作，则仍应传递一个掩码参数0xFFFFFFFF。简而言之，掩码参数和宽度参数大多是相互独立的。 - Robert Crovella

好的，谢谢。我现在觉得我已经完全理解了。我用shfl_sync替换了所有的shfl指令，我的代码仍然可以工作。顺便说一下，编译器似乎完全忽略非Volta架构的“mask”参数。我通过对所有shfl指令使用掩码“0x0”进行测试，发现它仍然可以工作：）。我的意思是这有点合理，因为一个warp中的所有线程都保证执行相同的指令。 - dari

你有没有想过，当你为sm_61编译时，为什么会从sm_30_intrinsics.hpp收到警告？ - Simon Huckett