c - 使用 CUDA 在本地内存中定义数组的可变大小

Question

是否有可能在设备函数中创建一个列表、数组、某些东西，列表/数组的大小在调用中作为参数……或者在调用时初始化的全局变量？

我希望这些列表中的一个能够工作:

unsigned int size1;

__device__ void function(int size2) {

    int list1[size1];
    int list2[size2];
}

有没有可能做一些聪明的事情来让这样的东西发挥作用？

Answer 1

有一种方法可以分配动态数量的共享内存——使用第三个启动内核参数:

__global__ void kernel (int * arr) 
{
    extern __shared__ int buf []; // size is not stated
    // copy data to shared mem:
    buf[threadIdx.x] = arr[blockIdx.x * blockDim.x + threadIdx.x];
    // . . . 
}
// . . . 
// launch kernel, set size of shared mem in bytes (k elements in buf):
kernel<<<grid, threads, k * sizeof(int)>>> (arr);

许多数组都有一个 hack:

__device__ void function(int * a, int * b, int k) // k elements in first list
{
    extern __shared__ int list1 [];
    extern __shared__ int list2 []; // list2 points to the same point as list1 does

    list1 [threadIdx.x] = a[blockIdx.x * blockDim.x + threadIdx.x];
    list2 [k + threadIdx.x] = b[blockIdx.x * blockDim.x + threadIdx.x];
    // . . .
}

您必须考虑:分配给所有 block 的内存。