c++ - 具有不同类型的CUDA每线程阵列

Question

我的CUDA内核的每个实例(即每个线程)都需要三个具有不同类型的私有(private)数组。

例如

__global__ void mykernel() {
    type1 a[aLen];
    type2 b[bLen];
    type3 c[cLen];

    ...
}

这些类型的大小在编译之前是未知的，并且 aLen， bLen和 cLen的长度是动态的。

当然，我必须为整个块分配一个共享内存实例。

void caller() {
    int threadsPerCUDABlock = ...
    int CUDABlocks = ...

    int threadMemSize = 
        aLen*sizeof(type1) + bLen*sizeof(type2) + cLen*sizeof(type3);

    int blockMemSize = threadsPerCUDABlock * threadMemSize;

    mykernel <<< CUDABlocks, threadsPerCUDABlock, blockMemSize >>>();
}

然后，每个线程的任务就是弄清楚共享内存的哪个分区是其私有(private)内存空间，以及如何将其划分为3种类型的子数组。在这种情况下，我将共享内存阵列组织为具有以下结构:

[ thread0_a, thread0_b, thread0_c,  thread1_a, ...]

我不确定如何最好地在内核中解压缩此结构。
我已经尝试传递每个线程的私有(private)空间的字节数，并且最初假设内存空间是1字节类型，例如 char:

mykernel <<< CUDABlocks, threadsPerCUDABlock, blockMemSize >>>(threadMemSize);

__global__ void mykernel(int threadMemSize) {

    extern __shared__ char sharedMem[];

    char* threadMem = &sharedMem[threadMemSize*threadIdx.x]
    type1 *a = (type1*) threadMem;
    type2 *b = (type2*) &a[aLen];
    type3 *c = (type3*) &b[bLen];

    ...
}

这没有用(尽管没有任何错误，但是很难调试)，但是我不相信它甚至在原理上也应该起作用。例如，我不能保证 type1， type2和 type3类型的大小严格减小。

那么总体上正确的范例是什么？就是说，用于解压缩不同类型和大小的多个每个线程数组的既定方法？

Answer 1

初步

通常，出于性能原因，人们对GPU计算很感兴趣-使他们的代码运行更快。因此，在尝试做出决策时，我们将以性能为指导。

我认为您在问题中提供的草图中的问题之一就是natural alignment requirement in CUDA之一。选择一个任意指针并将其类型转换为其他类型可能会违反此规定。如果您的代码中存在此类问题，则cuda-memcheck工具应该能够发现它。

将线程专用数组放入C++的典型位置是本地内存，我认为CUDA没什么不同。但是，CUDA C++至少不支持variable-length arrays。在您的问题中，您使用共享内存作为代理进行了草绘。您的想法(我假设)的含义之一是，尽管在编译时不知道这些数组的大小，但是必须有一个大小上限，因为共享内存可能会限制每个线程块低至48KB。因此，如果线程块中有1024个线程，则每个线程的最大组合数组大小将限制为48个字节。每个块有512个线程，可以想象每个线程有96个字节。如果使用共享内存，则可能是由于共享内存限制所致。

因此，另一种方法(如果您可以遵守这些下限)将是简单地上限所需的本地内存，并为每个线程静态定义该大小(或3)的本地内存数组。如已经提到的，必须将单个阵列划分在各个阵列之间，并注意对齐。但是考虑到您的方法建议的小尺寸(例如，总共约96个字节)，仅使用上限固定大小的本地数组(而不是共享内存)可能会很方便。

CUDA中的本地内存最终由与全局内存相同的物理资源(GPU DRAM内存)支持。然而，这种安排是这样的，如果每个线程正在访问其自己的本地存储器中的特定元素，那么如果该访问需要由DRAM服务，则线程间的影响将等同于合并访问。这意味着每线程本地存储以某种方式是交错的。如果出于性能考虑，如果我们提出自己的可变长度数组实现，那么这个交织特性也是我们要注意的。它同样适用于全局内存代理(以启用合并)或共享内存代理(以避免存储区冲突)。

除了出于性能原因而希望交织访问之外，而不是倾向于使用共享内存的可能的性能原因是共享内存的广泛使用可能会对占用率以及性能产生负面影响。该主题在其他许多地方都有介绍，因此在此不再赘述。

实现

本地记忆

如上所述，我认为关于使用共享内存的建议的一个隐含假设是，所需数组的实际大小必须有一些(合理较小)上限。如果是这种情况，最好使用3个分配了上限大小的数组:

const int Max_aLen = 9;
const int Max_bLen = 5;
const int Max_cLen = 9;
__global__ void mykernel() {
    type1 a[Max_aLen];
    type2 b[Max_bLen];
    type3 c[Max_cLen];

    ...
}

最多使用例如在我看来，用于本地内存的每个线程8 KB应该不是主要问题，但它可能取决于您的GPU和内存大小，并且 analysis mentioned/linked below应该指示任何问题。当然低水平/限制例如每个线程约96个字节应该不是问题。

全局记忆

我相信最简单，最灵活的方法是通过全局内存为指针和传递给内核的可变长度数组提供存储。这使我们可以通过例如 cudaMalloc，我们可以分别处理单独的数组，并且我们对对齐要求的关注相对较少。由于我们假装将这些全局数组当作线程专用来使用，因此我们将希望安排索引以创建每个线程的交错存储/访问，这将有助于合并。对于您的3数组示例，它可能看起来像这样:

#include <stdio.h>

typedef unsigned type1;
typedef char     type2;
typedef double   type3;

__global__ void mykernel(type1 *a, type2 *b, type3 *c) {

  size_t stride = (size_t)gridDim.x * blockDim.x;
  size_t idx = (size_t)blockIdx.x*blockDim.x+threadIdx.x;
  a[7*stride+idx] = 4;    // "local"  access to a
  b[0*stride+idx] = '0';  // "local"  access to b
  c[3*stride+idx] = 1.0;  // "local"  access to c
}

int main(){
  // 1D example
  type1 *d_a;
  type2 *d_b;
  type3 *d_c;
  // some arbitrary choices to be made at run-time
  size_t alen = 27;
  size_t blen = 55;
  size_t clen = 99;
  int nTPB = 256;
  int nBLK = 768;
  size_t grid = (size_t)nBLK*nTPB;
  // allocate
  cudaMalloc(&d_a, alen*grid*sizeof(type1));
  cudaMalloc(&d_b, blen*grid*sizeof(type2));
  cudaMalloc(&d_c, clen*grid*sizeof(type3));
  // launch
  mykernel<<<nBLK, nTPB>>>(d_a, d_b, d_c);
  cudaDeviceSynchronize();
}

对此方法的一种可能的批评是，它可能消耗的设备内存比本地内存方法要多(取决于 grid size relative to GPU type，它也可能消耗更少)。但是，这可以通过使用诸如 grid-stride looping之类的方法来限制网格大小来进行管理。

共享内存

由于我们只有一个指向共享内存的指针才能动态分配共享内存，因此，如果我们对共享内存执行某些操作，我们将必须特别注意对齐。这是分配和定位正确对齐的指针所需的计算类型的示例:

#include <stdio.h>

typedef unsigned type1;
typedef char     type2;
typedef double   type3;

__global__ void mykernel(int b_round_up, int c_round_up) {

  extern __shared__ char sdata[];
  type1 *a = (type1 *)sdata;
  type2 *b = (type2 *)(sdata + b_round_up);
  type3 *c = (type3 *)(sdata + c_round_up);
  size_t stride = blockDim.x;
  size_t idx = threadIdx.x;
  a[7*stride+idx] = 4;    // "local"  access to a
  b[0*stride+idx] = '0';  // "local"  access to b
  c[3*stride+idx] = 1.0;  // "local"  access to c
}

int main(){
  // 1D example
  // some arbitrary choices to be made at run-time
  int alen = 9;
  int blen = 5;
  int clen = 9;
  int nTPB = 256;
  int nBLK = 1;
  // calculate aligned shared mem offsets
  int b_round_up = (((nTPB*alen*sizeof(type1) + sizeof(type2)-1)/sizeof(type2))*sizeof(type2)); // round up
  int c_round_up = (((b_round_up + nTPB*blen*sizeof(type2) + sizeof(type3)-1)/sizeof(type3))*sizeof(type3)); // round up
  // allocate + launch
  mykernel<<<nBLK, nTPB, c_round_up + nTPB*clen*sizeof(type3)>>>(b_round_up,c_round_up);
  cudaDeviceSynchronize();
}

我并不是说我的任何代码都没有缺陷，但是从相对代码复杂性的 Angular 来看，本地或全局选项将是首选。我无法轻易想象共享内存实现将是首选的原因或情况。