在Cuda中使用最大共享内存
我不能使用超过48K的共享内存(在V100上，Cuda10.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。
如果我使用共享内存的<= 48K(例如，my_kernel<<<grd, blk, 48 * 1024, my_stream>>>)，它就能工作。
汇编标志是：
nvcc -std=c++14 -gencode arch=compute_70,code=sm_70 -Xptxas -v,-dlcm=cg
我遗漏了什么？
回答 1
Stack Overflow用户
发布于 2020-09-05 13:10:06
来自这里
计算能力7.x设备允许一个线程块来处理共享内存的全部容量: Volta为96 KB，图灵为64 KB。依赖于每个块超过48 KB的共享内存分配的内核是特定于体系结构的，因此它们必须使用动态共享内存(而不是静态大小的数组)，并且需要显式地选择使用cudaFuncSetAttribute()，如下所示：
cudaFuncSetAttribute(my_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 98304);当我将这一行添加到您显示的代码中时，无效的值错误就会消失。对于图灵设备，您可能希望将该数字从98304更改为65536。当然，65536对你的例子也是足够的，虽然不足以使用伏尔塔上的最大值，如问题标题中所述。
在安培上的类似时尚内核中，安培设备应该能够使用多达160 be的共享内存(cc 8.0)或100 be (cc 8.6)，使用上述选择机制动态分配，将数字98304更改为163840 (例如cc 8.0 )或102400 be( cc 8.6)。
请注意，以上内容涵盖了Volta (7.0)、图灵(7.5)和安培(8.x)两种情况。在7.x之前具有计算能力的GPU没有能力处理每个线程块超过48 to的地址。在某些情况下，这些GPU可能每个多处理器都有更多共享内存，但这是为了允许在某些线程块配置中占用更多内存。程序员没有能力使用超过48 no的每个线程块。
虽然它不涉及这里提供的代码(它已经在使用动态共享内存分配)，但摘录的文档引用中的注释指出，在支持它的设备上使用超过48 to的共享内存需要2件事：
- 上文已经描述的选择加入机制。
- 内核代码中的动态共享内存( 分配 )，而不是静态共享内存。
动态示例：
extern __shared__ int shared_mem[];静态示例：
__shared__ int shared_mem[1024];动态分配的共享内存还需要在内核启动配置参数中传递一个大小(问题中给出了一个示例)。
https://stackoverflow.com/questions/63757245
复制
腾讯云开发者
