在 CUDA 中，为什么我不能动态分配 2d 共享内存？答案

【问题标题】：In CUDA why can't I allocate 2d shared memory dynamically?在 CUDA 中，为什么我不能动态分配 2d 共享内存？
【发布时间】：2021-04-12 05:08:15
【问题描述】：

以下工作正常；

__extern__ float dyanimicSh1D[];

但以下不起作用：

__extern__ float dyanimicSh2D[][];

我想知道为什么会这样？

【问题讨论】：

这在 C++ 中的语法无效。为什么你认为它会在 CUDA 中工作？
但是在 C 中我可以做到：。 int two_d[10][20];
这不是一回事

标签： cuda gpu-shared-memory

【解决方案1】：

您不能这样做，因为编译器需要数组的宽度信息来生成执行正确索引的代码。

如果您像这样以静态方式分配共享内存：

__shared__ float sarr[24][12];

那么，您不仅要告诉分配/提供多少内存，还要提供数组的宽度（本例中为 12）。这一点很重要，因为这种类型的静态 2D 数组在底层并不被视为指针数组，而是一个平面分配，由编译器在编译时创建索引。

这样以后当你做这样的事情时：

float val = sarr[y][x];

编译器将获取sarr 指针，并进行指针运算以将x + (y*12) 添加到它，然后取消引用该指针以检索值。该计算中的 12 是在编译时发现的，并被编译器用于生成代码以进行索引。

做这样的事情：

extern __shared__ float sarr[][];

不向编译器提供数组宽度信息，因此无法生成编译时需要的索引，也是不允许的。

顺便说一句，这是可行的：

extern __shared__ float sarr[][12];

这是一个例子：

$ cat t46.cu
#include <cstdio>
__global__ void k(int x, int y){

  extern __shared__ float sarr[][12];
  for (int i = 0; i < 32; i ++)
    for (int j = 0; j < 12; j++)
      sarr[i][j] = i * 256 + j;
  float val = sarr[y][x];
  printf("%f\n", val);
}

int main(){

  k<<<1,1,128*12>>>(3,2);
  cudaDeviceSynchronize();
}
$ nvcc -o t46 t46.cu
$ cuda-memcheck ./t46
========= CUDA-MEMCHECK
515.000000
========= ERROR SUMMARY: 0 errors
$

【讨论】：