cuda - CUDA 中的空闲共享内存

Question

是否有任何应用程序级 API 可用于释放 CTA 在 CUDA 中分配的共享内存？我想将我的 CTA 重用于另一个任务，在开始该任务之前，我应该清理上一个任务使用的内存。

Answer 1

共享内存是在内核启动时静态分配的。您可以选择在内核中指定一个未分配大小的共享分配：

__global__ void MyKernel()
{
    __shared__ int fixedShared;
    extern __shared__ int extraShared[];
    ...
}

第三个内核启动参数然后指定有多少共享内存对应于未分配大小的分配。

MyKernel<<<blocks, threads, numInts*sizeof(int)>>>( ... );

为内核启动分配的共享内存总量是内核中声明的数量加上共享内存内核参数加上对齐开销的总和。你不能“释放”它——它在内核启动期间保持分配状态。

对于经历多个执行阶段并且需要将共享内存用于不同目的的内核，您可以做的是使用共享内存指针重用内存 - 在未调整大小的声明上使用指针算术。

就像是：

__global__ void MyKernel()
{
    __shared__ int fixedShared;
    extern __shared__ int extraShared[];
    ...
    __syncthreads();
    char *nowINeedChars = (char *) extraShared;
    ...
}

我不知道任何使用此习惯用法的 SDK 示例，尽管 threadFenceReduction 示例声明了 a__shared__ bool并且还使用共享内存来保存减少的部分总和。