cuda - 尝试使用超过 48 KB 的（静态）共享内存时出错

Question

我正在尝试__shared__使用sm_86. 如果变量的大小大于 48 KB，nvcc则报错并编译失败并报错

入口函数使用过多的共享数据

我正在使用 CUDA 11.2 运行 Visual Studio 2019 并已设置

配置属性 -> CUDA C/C++ -> 设备 -> 代码生成到 compute_86,sm_86

根据此处的表格，cc 8.6 的设备不应该允许每块 100 KB 的最大共享内存吗？

这是一个对我失败的简单代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#define SIZE 12288 /* x 4 byte = 48 KB <---- Works */
/* #define SIZE 12289 */ /* <---- Fails */

__global__ void kernel(int* d) {
    __shared__ int s[SIZE];
    s[threadIdx.x] = 1;
    d[threadIdx.x] = s[threadIdx.x];

}

int main() {
    int* d_a;
    cudaMalloc(&d_a, SIZE * sizeof(int));
    kernel <<<1, 1024 >>> (d_a);
    cudaDeviceSynchronize();
    cudaFree(d_a);
}

标签： cudashared-memory