cuda - CUDA 内核启动需要动态分配超过硬件容量的共享内存量

Question

据我了解，shared_mem_size下面代码中的变量是要为每个块分配的共享内存量。但是当总共享内存大小 ( BLOCKS_GRID * shared_mem_size) 超过硬件可用共享内存的总容量时会发生什么？

__global__ void KernelFunction ()
{
    extern __shared__ s_data[];
}

int main ()
{
    shared_mem_size = THREADS_BLOCK * sizeof(T);
    KernelFunction <<<BLOCK_GRID,THREADS_BLOCK,shared_mem_size>>> ();
}

Answer 1

在许多 CUDA 程序中，并非所有块都在同时执行。只要单个块请求的共享内存不超过 SM 上可用的内存（当前硬件为 16KB 或 48KB），那么至少一个块可以在该 SM 上执行。

假设我有一个由 1024 个块组成的网格。假设每个块需要 32KB 的共享内存。

在这种情况下，您可以在每个 SM 上驻留（即执行）一个线程块。如果我有一台有 4 个 SM 的机器，那么将执行 4 个线程块。随着这些线程块完成，可以启动更多线程块，直到消耗完所有 1024 个线程块。

如果单个线程块超过了可用的共享内存（例如，我请求 32KB，但 SM 配置为仅提供 16KB），那么我将收到内核启动错误/CUDA API 错误。