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共享内存数组变量)有一个类似的问题,他们被告知创建一个单独的数组,作为两个数组的数组,只需调整索引以正确匹配数组.虽然这似乎做了我想要的,但它看起来非常混乱.有没有办法解决这个问题,以便我仍然可以维护两个独立的数组,每个数组的大小都被用户定义为输入？

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
$

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