c - 为什么我的内核没有超过共享内存限制？

Question

我正在从 matlab 调用 CUDA 内核。

我之前被告知(David Kirk 的书)每个线程只能占用 16kb 的共享内存，但我能够消耗的远不止于此:

__global__ void plain(float* arg)
{

    __shared__ float array[12000];
    int k;

    for (k=1;k<12000; k++)
    {
        array[k]=1;
    }   
}

CUDA C 报告 float 为 4 个字节，这意味着总数组大小为 48Kb，大于 12Kb。运行正常，怎么会这样？

我也被告知 GPU shared memory size is very small - what can I do about it?每个 block 的最大共享内存很重要。我卡的每个 block 的最大共享内存是49152 字节，但我能够以每个 block 1000 个线程运行上述代码。

似乎每个 block 将使用 49Kb，这是不对的。是不是SM一次只服务一个 block ，in dong保留了每个线程 block 只能有49Kb的条件？

每个 block 的 49Kb 共享内存如何与每个线程的 16Kb 共享内存相协调？

谢谢

Answer 1

共享内存按线程 block 分配，每个 SM 有多达 48 KB 的可用计算能力 2.0 及更高版本。因此，在给定的 SM 上，您可以运行一个占用整个 48 KB 的线程 block ，或者说，三个线程 block ，每个线程 block 分配 16 KB。每个 SM 16 KB 共享内存的限制适用于 < 2.0 的计算能力。与按线程 block 分配的共享内存相反，按线程分配本地内存(“本地”表示“线程本地”)。