微信公众号搜"智元新知"关注
微信扫一扫可直接关注哦!

在 CUDA 中为什么不能动态分配二维共享内存?

如何解决在 CUDA 中为什么不能动态分配二维共享内存?

以下工作正常;

__extern__ float dyanimicSh1D[];

但以下方法不起作用:

__extern__ float dyanimicSh2D[][];

我想了解为什么会这样?

解决方法

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

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

__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
$

版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。