CUDA 静态共享内存释放

【问题标题】：CUDA static shared memory deallocationCUDA 静态共享内存释放
【发布时间】：2013-09-17 08:00:39
【问题描述】：

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

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

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

【问题讨论】：

在 C/C++ 中，无法释放静态定义的数组。为什么不从内核启动时动态分配最坏情况所需的共享内存量（extern __shared__ ... 然后myKernel<<<numBlocks, numThreads, sh_mem_size>>>）？

标签： cuda gpgpu

【解决方案1】：

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

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

extern __shared__ float foo[];

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

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

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

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

通过同样的概念，你可以动态改变你正在处理的数组的大小。

【讨论】：