CUDA 的变形改组

Question

我需要做一个像这样的扭曲洗牌：

在此图片上，线程数限制为 8 以使其可读。如果我阅读了 Nvidia SDK 和 ptx 手册，shuffle 指令应该可以完成这项工作，特别是 shfl.idx.b32 d[|p], a, b, c; ptx 指令。

从 manual 我读到：

Each thread in the currently executing warp will compute a source lane
index j based on input operands b and c and the mode. If the computed
source lane index j is in range, the thread will copy the input operand
a from lane j into its own destination register d;

因此，提供 b 和 c 的适当值，我应该能够通过编写这样的函数来完成（灵感来自 CUDA SDK __shufl 原始实现）。

  __forceinline__ __device __ float shuffle(float var){
   float ret;
   int srcLane = ???
   int c = ???
   asm volatile ("shfl.idx.b32 %0, %1, %2, %3;" : "=f"(ret) : "f"(var), "r"(srcLane), "r"(c));
  return ret;

}

如果可能的话，srcLane和c的常数是多少？我无法确定它们（我使用的是 CUDA 8.0）。

最佳，

蒂莫咖啡厅

Answer 1

我建议使用 CUDA intrinsic 而不是 PTX（或内联 ASM）来执行此操作。然而，以下代码演示了这两种方法：

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

__global__ void k(){

  int i = threadIdx.x;
  int j = i;
  if (i<4) j*=2;
  if ((i>3) && (i<8)) j-=(7-i);
  int k = __shfl_sync(0x0FFU, i+100, j);
  printf("lane: %d, result: %d\n", i, k);
}
__forceinline__ __device__ float shuffle(float var, int lane){
   float ret;
   int srcLane = lane;
   int c = 0x1F;
   asm volatile ("shfl.idx.b32 %0, %1, %2, %3;" : "=f"(ret) : "f"(var), "r"(srcLane), "r"(c));
  return ret;
}

__global__ void k1(){

  int i = threadIdx.x;
  int j = i;
  if (i<4) j*=2;
  if ((i>3) && (i<8)) j-=(7-i);
  float k = shuffle((float)(i+100), j);
  printf("lane: %d, result: %f\n", i, k);
}


int main(){


  k<<<1,8>>>();
  cudaDeviceSynchronize();
  k1<<<1,8>>>();
  cudaDeviceSynchronize();
}
$ nvcc -arch=sm_35 -o t54 t54.cu
$ cuda-memcheck ./t54
========= CUDA-MEMCHECK
lane: 0, result: 100
lane: 1, result: 102
lane: 2, result: 104
lane: 3, result: 106
lane: 4, result: 101
lane: 5, result: 103
lane: 6, result: 105
lane: 7, result: 107
lane: 0, result: 100.000000
lane: 1, result: 102.000000
lane: 2, result: 104.000000
lane: 3, result: 106.000000
lane: 4, result: 101.000000
lane: 5, result: 103.000000
lane: 6, result: 105.000000
lane: 7, result: 107.000000
========= ERROR SUMMARY: 0 errors
$

使用 CUDA 内在函数（第一种方法），唯一真正的任务是计算源通道索引。根据您的模式，我编写了一些代码来执行此操作并将其放入变量 j.

Answer 2

您在 shuffle 操作中尝试做的是能够动态索引 shuffle 运行的源通道。人们需要了解 shuffle 命令的任何变体 (__shfl, __shfl_up, __shfl_down, __shfl_xor) 的第二个参数都需要一个常量值，并且此参数对于一个区域内的所有车道都是相同经。您可以通过指定 width 在 warp 中对线程进行分组。因此，例如，通过指定

float var = ...
__shfl_xor(var, 3, 4);

车道排列如下：

0 1 2 3
   |
3 2 1 0

因此，要回答您的问题，不可能通过任何类型的单个 __shuffle 操作来完成。但是您可以通过组合多个具有不同第二参数的 __shuffle 命令来实现它。

Answer 3

罗伯特已经圆满地回答了这个问题。我已经实现了下面的代码，显示了完整扭曲的排列。

#include <stdio.h>

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort = true)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) { getchar(); exit(code); }
    }
}

__global__ void shufflingKernel(double *d_data, double *d_result, int *d_perm){

    unsigned mask = __activemask(); 
    int tid = threadIdx.x;
    int srcLane = d_perm[tid];
    double var = d_data[tid];
    //d_result[tid] = __shfl_sync(0xFFFFFFFF, var, srcLane);
    d_result[tid] = __shfl_sync(mask, var, srcLane);
}

int main(){

    const int N = 32;

    double h_data[32] = { 3.4, 42.2, 2., -1., 10., 11., 2., -1., 10., 33., 2.3, 11., 44., 0., -33., -21.,
        4.4, 43.2, 3., -2., 13., 15., 222., -90., 17., 30., 11.3, 7., 22., 100., -30., -91. };
    double *h_result = (double *)malloc(N * sizeof(double));
    int h_perm[32] = { 6, 11, 9, 2, 5, 23, 31, 0, 3, 27, 29, 1, 28, 30, 17, 13, 10, 8, 4, 22, 7, 18, 24, 12, 20,
        19, 16, 26, 21, 15, 25, 14 };

    int *d_perm; gpuErrchk(cudaMalloc(&d_perm, N * sizeof(int)));
    double *d_data; gpuErrchk(cudaMalloc(&d_data, N * sizeof(double)));
    double *d_result; gpuErrchk(cudaMalloc(&d_result, N * sizeof(double)));
    gpuErrchk(cudaMemcpy(d_perm, &h_perm[0], N * sizeof(int), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_data, &h_data[0], N * sizeof(double), cudaMemcpyHostToDevice));

    shufflingKernel << <1, 32>> >(d_data, d_result, d_perm);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaMemcpy(h_result, d_result, N * sizeof(double), cudaMemcpyDeviceToHost));

    for (int k = 0; k < N; k++) {
        printf("k = %d; Original = %f; New = %f; Check = %f\n", k, h_data[k], h_result[k], h_data[h_perm[k]]);
    }

}

请注意，不是使用 0xFFFFFFFF 作为活动线程的掩码，而是使用意义上的 warp-level 原语 __activemask() 更安全。

CUDA 的变形改组

Warp shuffling for CUDA

cuda

shuffle

ptx

gpu-warp