在 Cuda 中使用最大共享内存答案

【问题标题】：Using maximum shared memory in Cuda在 Cuda 中使用最大共享内存
【发布时间】：2021-06-16 06:00:54
【问题描述】：

我无法使用超过 48K 的共享内存（在 V100、Cuda 10.2 上）

我打电话

cudaFuncSetAttribute(my_kernel,
                     cudaFuncAttributePreferredSharedMemoryCarveout,
                     cudaSharedmemCarveoutMaxShared);

在第一次启动my_kernel 之前。

我使用启动边界和my_kernel内部的动态共享内存：

__global__
void __launch_bounds__(768, 1)
my_kernel(...)
{
    extern __shared__ float2 sh[];
    ...
}

内核是这样调用的：

dim3 blk(32, 24); // 768 threads as in launch_bounds.

my_kernel<<<grd, blk, 64 * 1024, my_stream>>>( ... );

cudaGetLastError() 内核调用后返回cudaErrorInvalidValue。

如果我使用 my_kernel<<<grd, blk, 48 * 1024, my_stream>>>），它可以工作。

编译标志是：

nvcc -std=c++14 -gencode arch=compute_70,code=sm_70 -Xptxas -v,-dlcm=cg

我错过了什么？

【问题讨论】：

标签： cuda

【解决方案1】：

来自here：

计算能力 7.x 设备允许单个线程块来处理共享内存的全部容量：Volta 为 96 KB，Turing 为 64 KB。依赖于每个块超过 48 KB 的共享内存分配的内核是特定于架构的，因此它们必须使用动态共享内存（而不是静态大小的数组）并且需要使用 cudaFuncSetAttribute 显式选择加入()如下：

cudaFuncSetAttribute(my_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 98304);

当我将该行添加到您显示的代码中时，无效值错误就会消失。对于图灵设备，您可能希望将该数字从 98304 更改为 65536。当然 65536 对于您的示例也足够了，但不足以使用 volta 上可用的最大值，如问题标题中所述。

在 Ampere 设备上的 similar fashion 内核中，应该能够使用高达 160KB 的共享内存 (cc 8.0) 或 100KB (cc 8.6)，使用上述选择加入机制动态分配，编号 98304 已更改到 163840（例如，对于 cc 8.0）或 102400（对于 cc 8.6）。

请注意，以上内容涵盖了 Volta (7.0) Turing (7.5) 和 Ampere (8.x) 的情况。在 7.x 之前具有计算能力的 GPU 无法处理每个线程块超过 48KB 的容量。在某些情况下，这些 GPU 的每个多处理器可能有更多的共享内存，但这是为了在某些线程块配置中允许更大的占用。程序员无法使用每个线程块超过 48KB。

虽然它与此处提供的代码无关（它已经使用动态共享内存分配），但请注意摘录文档引用中的内容，即在支持它的设备上使用超过 48KB 的共享内存需要两件事：

上面已经描述了选择加入机制
内核代码中的动态而不是静态共享内存allocation。

动态示例：

extern __shared__ int shared_mem[];

静态示例：

__shared__ int shared_mem[1024];

动态分配的共享内存还需要在内核启动配置参数中传递一个大小（问题中给出了一个示例）。

【讨论】：