cuda - CUDA 中的动态共享内存

Question

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

现在我可以按照类似于手册中的示例( http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared )，例如

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 )有一个类似的问题，他们被告知创建一个单一的数组，作为两个数组的数组，并简单地调整索引以正确匹配数组。虽然这似乎确实符合我的要求，但它看起来非常困惑。有什么办法可以解决这个问题，以便我仍然可以维护两个独立的数组，每个数组的大小都由用户定义为输入？

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
$

如果您有混合数据类型数组的随机排列，您需要手动对齐数组起点(并请求足够的共享内存)或使用对齐指令(并确保请求足够的共享内存)，或使用有助于对齐的结构。