预赛
通常，人们出于性能原因对 GPU 计算感兴趣 - 以使他们的代码运行得更快。因此，我们会在尝试做出决定时将性能作为指导。
我认为您在问题中提供的草图会遇到的问题之一是natural alignment requirement in CUDA。选择任意指针并将其类型转换为不同的类型可能会与此冲突。如果您的代码中有这样的问题，cuda-memcheck 工具应该能够发现它。
在 C++ 中放置线程私有数组的典型位置是 local 内存，我认为 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];
...
}
最多使用例如在我看来，本地内存每个线程 8kbytes 不应该是一个主要问题，但这可能取决于您的 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();
}
我并不是说我的任何代码都没有缺陷，但是从相对代码复杂性的角度来看，您可以看到本地或全局选项是首选。我无法轻易想象共享内存实现更受欢迎的原因或情况。