我无法使用超过 48K 的共享内存(在 V100、Cuda 10.2 上)
我打电话
cudaFuncSetAttribute(my_kernel,
cudaFuncAttributePreferredSharedMemoryCarveout,
cudaSharedmemCarveoutMaxShared);
Run Code Online (Sandbox Code Playgroud)
在my_kernel第一次启动之前。
我在内部使用启动边界和动态共享内存my_kernel:
__global__
void __launch_bounds__(768, 1)
my_kernel(...)
{
extern __shared__ float2 sh[];
...
}
Run Code Online (Sandbox Code Playgroud)
内核是这样调用的:
dim3 blk(32, 24); // 768 threads as in launch_bounds.
my_kernel<<<grd, blk, 64 * 1024, my_stream>>>( ... );
Run Code Online (Sandbox Code Playgroud)
cudaGetLastError()内核调用返回后cudaErrorInvalidValue。
如果我使用 <= 48 K 的共享内存(例如,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