CUDA 9 шфл против шфл_синк

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

Answer 1

Если вы прочитаете руководство по программированию CUDA 9RC (раздел B.15), установленное вместе с вашей копией CUDA 9RC, вы увидите, что новый __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 принимают маску, указывающую потоки, участвующие в вызове. Бит, представляющий идентификатор полосы потока, должен быть установлен для каждого участвующего потока, чтобы гарантировать, что они должным образом сходятся до того, как аппаратное обеспечение выполнит встроенную функцию. Все неперешедшие потоки, названные в маске, должны выполнять одну и ту же функцию с одинаковой маской, иначе результат не определен.

Поэтому, если мы изменим ваш код для соответствия этому, мы получим ожидаемый результат:

$ 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
$