1

我需要做一个看起来像这样的经纱改组： 经线改组

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

从我读到的手册：

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;

因此，提供适当的band值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）。

最好的，

蒂莫咖啡

4

3 回答 3

3

我建议使用CUDA 内在函数而不是 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。

于 2018-03-10T03:32:52.683 回答
1

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

#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]]);
    }

}

请注意，在 CUDA 中的 Shuffle 指令不工作的意义上，0xFFFFFFFF使用 warp 级原语而不是使用活动线程的掩码更安全。__activemask()

于 2018-03-15T18:03:28.790 回答
-2

您在shuffle操作中尝试做的是能够动态索引 shuffle 操作的源通道。人们需要了解shuffle命令 ( __shfl, __shfl_up, __shfl_down, __shfl_xor) 的任何变化都需要其第二个参数的恒定值，并且该参数对于经线中的所有通道都是相同的。您可以通过指定width. 因此，例如，通过指定

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

车道排列将如下所示：

0 1 2 3
   |
3 2 1 0

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

于 2018-03-11T17:08:00.107 回答