CUDA中的动态共享内存

Question

有类似于我将要询问的问题，但我觉得它们都没有涉及到我真正寻找的核心。我现在拥有一个需要将两个数组定义到共享内存中的CUDA方法。现在，数组的大小由一个在程序执行开始后读入程序的变量给出。因此，我不能使用该变量来定义数组的大小，因为定义共享数组的大小需要在编译时知道值。我不想像这样做 __shared__ double arr1[1000] ，因为手动输入大小对我而言是无用的，因为它将根据输入而改变。同样地，我也不能使用 #define 来创建一个常量大小。

现在，我可以遵循类似于manual中的示例。

extern __shared__ float array[];
__device__ void func()      // __device__ or __global__ function
{
    short* array0 = (short*)array; 
    float* array1 = (float*)&array0[128];
    int*   array2 =   (int*)&array1[64];
}

但是这仍然存在一个问题。根据我所读的，定义共享数组总是使内存地址成为第一个元素。这意味着我需要使我的第二个数组向右移动第一个数组的大小，就像在这个例子中所做的那样。但是第一个数组的大小取决于用户输入。

另一个问题（Cuda Shared Memory array variable）有类似的问题，他们被告知创建一个单一的数组，作为两个数组的数组，并简单地调整索引以正确匹配数组。虽然这似乎可以实现我想要的效果，但看起来非常混乱。有没有办法解决这个问题，以便我仍然可以维护两个独立的数组，每个数组的大小都由用户输入定义？

- zephyr

Answer 1

使用CUDA动态共享内存时，只有一个指针被传递给内核，该指针定义了以字节为单位的请求/分配区域的起始位置：

extern __shared__ char array[];

没有其他处理的方法。但是这并不妨碍您拥有两个用户大小的数组。以下是一个实例：

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

__global__ void my_kernel(unsigned arr1_sz, unsigned arr2_sz){

  extern __shared__ char array[];

  double *my_ddata = (double *)array;
  char *my_cdata = arr1_sz*sizeof(double) + array;

  for (int i = 0; i < arr1_sz; i++) my_ddata[i] = (double) i*1.1f;
  for (int i = 0; i < arr2_sz; i++) my_cdata[i] = (char) i;

  printf("at offset %d, arr1: %lf, arr2: %d\n", 10, my_ddata[10], (int)my_cdata[10]);
}

int main(){
  unsigned double_array_size = 256;
  unsigned char_array_size = 128;
  unsigned shared_mem_size = (double_array_size*sizeof(double)) + (char_array_size*sizeof(char));
  my_kernel<<<1,1, shared_mem_size>>>(256, 128);
  cudaDeviceSynchronize();
  return 0;
}


$ nvcc -arch=sm_20 -o t501 t501.cu
$ cuda-memcheck ./t501
========= CUDA-MEMCHECK
at offset 10, arr1: 11.000000, arr2: 10
========= ERROR SUMMARY: 0 errors
$

如果你有一个随机排列的混合数据类型数组，你需要手动对齐数组起始点（并请求足够的共享内存），或者使用对齐指令（并确保请求足够的共享内存），或者使用结构体来帮助对齐。

- Robert Crovella

---

谢谢。我有点想到它可以这样工作，但我还是想确认一下。快速问题，如果我使用 char *my_cdata = (char*)&my_ddata[arr1_sz];，它仍然有效吗？ - zephyr

这现在只是关于C语言的问题，对吧？是的，我认为它是有效的。 - Robert Crovella