寄存器中 AVX 置换和随机播放的 CUDA 翻译答案

【问题标题】：CUDA translation of AVX permute and shuffle in registers寄存器中 AVX 置换和随机播放的 CUDA 翻译
【发布时间】：2019-12-22 03:26:10
【问题描述】：

我正在尝试将 AVX 例程转换为 CUDA，并且大部分工作都非常简单。然而，由于缺乏简单的例子，我无法理解这个翻译的两段。

如何对寄存器浮点变量（长度始终为 32）执行任意排列？我已经看到 __shfl_sync 会执行此操作的建议，但没有显示此示例的示例。我想用长度为 8 的数组做的一个简单案例的 numpy 版本：
```
"""
a == some float32 array of length 8;
specific  patterns will always cycle mod 4
"""
b = a[[3,2,1,0,7,6,5,4]] 
```

如何将两个寄存器浮点数合并为一个寄存器浮点数？在 numpy 中，一个简单的例子是：

"""
a == some float32 array of length 8 
b == some other float32 array of length 8
specific  patterns will always cycle mod 4 
"""
c = numpy.array([a[0],a[1], b[0],b[1], 
                 a[4],a[5], b[4],b[5]])

对于任何了解 AVX 内在函数的人，问题 1 与 _mm256_permute_ps 的翻译有关，问题 2 与 _mm256_shuffle_ps 的翻译有关。

【问题讨论】：

真正简短的回答是 CUDA 中不存在范式。 SIMD 指令的数量非常少，但它们和所有寄存器一样只有 32 位宽
case 1 对我来说看起来像是一个简单的 __shfl_sync 应用程序。找不到示例的建议很奇怪。案例2似乎也很简单。可能我对 AVX 还不够了解。
我可以找到大量具有特殊轮班或广播的示例，但对于经纱中的通用重新排序却一无所获。如果它是“...一个简单的应用程序 __shfl_sync...”，那真是个好消息，但我绝对会欣赏一个示例或链接。仅供参考，我要翻译的 SIMD 例程是用于快速批量 4x4 矩阵求逆，我想与您在 stackoverflow.com/questions/55007384/… 上发布的进行比较。你的例子很有教育意义，我非常感谢他们。
好吧，我大概知道 talonmies 正在做什么。您想在单个 CUDA 线程的范围内执行此操作，对吗？如果是这样，那么不，我无法帮助你。没有包含 32 个float 数量的 CUDA 寄存器或寄存器结构。如果你想将 AVX 翻译成 CUDA，一次使用 32 个线程可能会更有效率。
很抱歉，我没有指定假定的经纱大小为 32。我认为我想做的事情是可能的，正如 www2.maths.ox.ac.uk/~gilesm/cuda/2019/lecture_04.pdf 中的第 4/5 页所建议的那样

标签： cuda simd avx

【解决方案1】：

如何对寄存器浮点变量（长度始终为 32）执行任意排列？我已经看到 __shfl_sync 会执行此操作的建议，但没有显示此示例的示例。我想用长度为 8 的数组做的一个简单案例的 numpy 版本：

a == 一些长度为 8 的 float32 数组；特定模式将始终循环 mod 4 """ b = a[[3,2,1,0,7,6,5,4]]

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

__global__ void k(int *pattern){

  float my_val = (float)threadIdx.x + 0.1f;
  my_val = __shfl_sync(0xFFFFFFFF, my_val, pattern[threadIdx.x]);
  printf("warp lane: %d, val: %f\n", threadIdx.x&31, my_val);
}

int main(){

  int pattern[32] = {3,2,1,0,7,6,5,4};
  for (int i = 8; i<32; i++) pattern[i] = i;
  int *d_pattern;
  cudaMalloc(&d_pattern, sizeof(pattern));
  cudaMemcpy(d_pattern, pattern, sizeof(pattern), cudaMemcpyHostToDevice);
  k<<<1,32>>>(d_pattern);
  cudaDeviceSynchronize();
}


$ nvcc -o t1486 t1486.cu
$ cuda-memcheck ./t1486
========= CUDA-MEMCHECK
warp lane: 0, val: 3.100000
warp lane: 1, val: 2.100000
warp lane: 2, val: 1.100000
warp lane: 3, val: 0.100000
warp lane: 4, val: 7.100000
warp lane: 5, val: 6.100000
warp lane: 6, val: 5.100000
warp lane: 7, val: 4.100000
warp lane: 8, val: 8.100000
warp lane: 9, val: 9.100000
warp lane: 10, val: 10.100000
warp lane: 11, val: 11.100000
warp lane: 12, val: 12.100000
warp lane: 13, val: 13.100000
warp lane: 14, val: 14.100000
warp lane: 15, val: 15.100000
warp lane: 16, val: 16.100000
warp lane: 17, val: 17.100000
warp lane: 18, val: 18.100000
warp lane: 19, val: 19.100000
warp lane: 20, val: 20.100000
warp lane: 21, val: 21.100000
warp lane: 22, val: 22.100000
warp lane: 23, val: 23.100000
warp lane: 24, val: 24.100000
warp lane: 25, val: 25.100000
warp lane: 26, val: 26.100000
warp lane: 27, val: 27.100000
warp lane: 28, val: 28.100000
warp lane: 29, val: 29.100000
warp lane: 30, val: 30.100000
warp lane: 31, val: 31.100000
========= ERROR SUMMARY: 0 errors
$

对于问题 2，我唯一能想到的似乎微不足道。正如我对问题 1 的回答中所建议的那样，考虑 32 项 float 数组的一种方法是让数组“散布”在一个扭曲上。我相信这与 AVX 样式处理最对应。

如果我们遵循这一点，那么问题 2 的代码可能很简单：

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

__global__ void k(int *pattern){

  float my_vals[2] = {1.1f, 2.2f};
  float my_val = my_vals[pattern[threadIdx.x]];
  printf("warp lane: %d, val: %f\n", threadIdx.x&31, my_val);
}

int main(){

  int pattern[32] = {0,0,1,1,0,0,1,1};
  for (int i = 8; i<32; i++) pattern[i] = 0;
  int *d_pattern;
  cudaMalloc(&d_pattern, sizeof(pattern));
  cudaMemcpy(d_pattern, pattern, sizeof(pattern), cudaMemcpyHostToDevice);
  k<<<1,32>>>(d_pattern);
  cudaDeviceSynchronize();
}


$ nvcc -o t1487 t1487.cu
$ cuda-memcheck ./t1487
========= CUDA-MEMCHECK
warp lane: 0, val: 1.100000
warp lane: 1, val: 1.100000
warp lane: 2, val: 2.200000
warp lane: 3, val: 2.200000
warp lane: 4, val: 1.100000
warp lane: 5, val: 1.100000
warp lane: 6, val: 2.200000
warp lane: 7, val: 2.200000
warp lane: 8, val: 1.100000
warp lane: 9, val: 1.100000
warp lane: 10, val: 1.100000
warp lane: 11, val: 1.100000
warp lane: 12, val: 1.100000
warp lane: 13, val: 1.100000
warp lane: 14, val: 1.100000
warp lane: 15, val: 1.100000
warp lane: 16, val: 1.100000
warp lane: 17, val: 1.100000
warp lane: 18, val: 1.100000
warp lane: 19, val: 1.100000
warp lane: 20, val: 1.100000
warp lane: 21, val: 1.100000
warp lane: 22, val: 1.100000
warp lane: 23, val: 1.100000
warp lane: 24, val: 1.100000
warp lane: 25, val: 1.100000
warp lane: 26, val: 1.100000
warp lane: 27, val: 1.100000
warp lane: 28, val: 1.100000
warp lane: 29, val: 1.100000
warp lane: 30, val: 1.100000
warp lane: 31, val: 1.100000
========= ERROR SUMMARY: 0 errors
$

如果这是一个学习练习，那就太好了。如果您的兴趣是对 4x4 批处理矩阵求逆进行稳健的实现，我建议您使用 CUBLAS。

【讨论】：

您的批处理 4x4 逆 (t411.cu) 版本在我的旧 Quadro K4200 上的性能比 CUBLAS 高 9 倍，而且它的性能也略胜于 MAGMA。

【解决方案2】：

对于问题 2，我有第二个解决方案，在 Robert 发布他的问题之前我已经解决了这个问题。我将不得不多研究一下接受的内容，但在这一点上，我很高兴有多种选择。

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

__device__ unsigned pat[4];
const unsigned hpat[4] = {1, 1, 0, 0};

__global__ void k(int *pattern){

  float my_val = (float)threadIdx.x + 0.0f;
  float my_val1 = (float)threadIdx.x + 32.0f;
  float out_val = 0.0;
  out_val = my_val*pat[threadIdx.x%4];
  out_val += __shfl_up_sync(0xFFFFFFFF, my_val1, 2, 4)*(1-pat[threadIdx.x%4]);
  printf("warp lane: %d, val: %f\n", threadIdx.x&31, out_val);
}

int main(){

  int pattern[32] = {3,2,1,0,7,6,5,4};
  for (int i = 8; i<32; i++) pattern[i] = i;
  int *d_pattern;
  cudaMemcpyToSymbol(pat, hpat, 4*sizeof(unsigned));
  cudaMalloc(&d_pattern, sizeof(pattern));
  cudaMemcpy(d_pattern, pattern, sizeof(pattern), cudaMemcpyHostToDevice);
  k<<<1,32>>>(d_pattern);
  cudaDeviceSynchronize();
}

$ nvcc -o t1486 t1486.cu
$ ./t1486
warp lane: 0, val: 0.000000
warp lane: 1, val: 1.000000
warp lane: 2, val: 32.000000
warp lane: 3, val: 33.000000
warp lane: 4, val: 4.000000
warp lane: 5, val: 5.000000
warp lane: 6, val: 36.000000
warp lane: 7, val: 37.000000
warp lane: 8, val: 8.000000
warp lane: 9, val: 9.000000
warp lane: 10, val: 40.000000
warp lane: 11, val: 41.000000
warp lane: 12, val: 12.000000
warp lane: 13, val: 13.000000
warp lane: 14, val: 44.000000
warp lane: 15, val: 45.000000
warp lane: 16, val: 16.000000
warp lane: 17, val: 17.000000
warp lane: 18, val: 48.000000
warp lane: 19, val: 49.000000
warp lane: 20, val: 20.000000
warp lane: 21, val: 21.000000
warp lane: 22, val: 52.000000
warp lane: 23, val: 53.000000
warp lane: 24, val: 24.000000
warp lane: 25, val: 25.000000
warp lane: 26, val: 56.000000
warp lane: 27, val: 57.000000
warp lane: 28, val: 28.000000
warp lane: 29, val: 29.000000
warp lane: 30, val: 60.000000
warp lane: 31, val: 61.000000

【讨论】：