在Cuda中使用最大共享内存

问题描述

我不能使用超过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,48 * 1024,my_stream>>>），它将起作用。

编译标志为：

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

我想念什么？

解决方法

来自here：

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

cudaFuncSetAttribute(my_kernel,cudaFuncAttributeMaxDynamicSharedMemorySize,98304);

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

在similar fashion中，安培设备上的内核应该能够使用上述选择加入机制动态分配多达160KB的共享内存，并且编号98304更改为163840。

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