为什么我的内核不超过共享内存限制？

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报告浮点数为4个字节,这意味着总数组大小为48Kb,大于12Kb.它运行良好,那怎么可能呢？

我也被告知在 GPU中共享内存的大小非常小 - 我该怎么办呢？ 每个块的最大共享内存很重要.我的卡的每块最大共享内存是49152字节,但我能够运行上面的代码,每块1000个线程.

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

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

谢谢

Answer 1

每个线程块分配共享内存,每个SM可用多达48 KB,计算能力为2.0及以上.因此,在给定的SM上,您可以运行单个线程块,该块占用整个48 KB,或者说三个线程块,每个块分配16 KB.每个SM的16 KB共享内存限制适用于<2.0的计算能力.与每个线程块分配的共享存储器相反,每个线程分配本地存储器("本地",意思是"线程本地").