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

【问题标题】：How come my kernel doesn't exceed the shared memory limit?为什么我的内核没有超过共享内存限制？
【发布时间】：2012-08-22 19:56:26
【问题描述】：

我正在从 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 报告浮点数为 4 个字节，这意味着总数组大小为 48Kb，大于 12Kb。运行正常，怎么会这样？

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

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

每块 49Kb 共享内存与每线程 16Kb 共享内存如何协调？

谢谢

【问题讨论】：

共享内存是按线程块分配的，每个 SM 有多达 48 KB 的可用空间，计算能力为 2.0 及更高版本。因此，在给定的 SM 上，您可能正在运行一个消耗整个 48 KB 的线程块，或者说，三个线程块每个分配 16 KB。每个 SM 16 KB 共享内存的限制适用于计算能力
@njuffa，您的评论应该是一个答案。 :)
@harrism：谢谢，根据您的建议将评论复制到答案中。

【解决方案1】：

共享内存是按线程块分配的，每个 SM 最多可使用 48 KB，计算能力为 2.0 及更高版本。因此，在给定的 SM 上，您可能正在运行一个消耗整个 48 KB 的线程块，或者说，三个线程块每个分配 16 KB。每个 SM 16 KB 共享内存的限制适用于计算能力

【讨论】：

【解决方案2】：

线程没有共享内存。您的代码使用“块”共享内存（CUDA 中没有其他共享内存）

【讨论】：