cuda - CUDA 静态共享内存释放

Question

有什么方法可以释放在同一个 CUDA 内核中预先分配的共享内存？例如，在内核内部，我曾经定义过

__shared__ unsigned char flag;
__shared__ unsigned int values [ BLOCK_DIM ];

稍后在代码内部，我需要定义一个数组，考虑到先前定义的共享内存超过为块设置的共享内存限制。如果没有重复使用先前定义的共享内存的肮脏工作，我怎么能做到这一点？或者 NVCC 足够聪明，可以识别内核跟踪中的依赖关系，并在使用共享变量完成时释放它？我的设备是 GeForce GTX 780 (CC=3.5)。

Answer 1

在 C/C++ 中，无法释放静态定义的数组。

您可能希望动态分配最坏情况所需的共享内存量，如下所示。添加

extern __shared__ float foo[];

在内核函数中并启动你的内核函数

myKernel<<<numBlocks, numThreads, sh_mem_size>>> (...);  

请记住，您可以通过使用指针来管理多个数组。查看CUDA C 编程指南了解更多详细信息。例如，引用指南

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];
}

通过相同的概念，您可以动态更改正在处理的数组的大小。