cuda shared memory 静态分配和动态分配_maxdynamicsharedmemoryperblocksize-CSDN博客

本文深入探讨CUDA编程中共享内存的使用，包括静态和动态分配的方法，以及它们在GPU内存布局中的作用。通过实例展示了不同分配方式下，共享内存的地址空间特点及限制。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

静态分配

加上前缀 shared

__shared__ int _ss[1024];1

动态分配

当我们在编程时，不清楚shared memory 数组开多大，就要用到动态分配。
分为两部分：
1，声明

extern __shared__ int _s[];

2, 在调用kernel 时加上数组的大小。

xxx_kernel<<<grid, block, sharedMemSize>>>();1

内存分布

下面通过一个例子，说明同时使用静态和动态分配时，内存分配情况。
kernel 代码：

__global__ void sharedMemTest()
{
    extern __shared__ int _s[];
    __shared__ int _ss[1024];
    if (threadIdx.x==0)
        printf("blockIdx.x is %d s is at %x, ss is at %x\n", blockIdx.x, _s, _ss);
}

调用kernel代码：

{
        dim3 block(32);
        dim3 grid(32);
        sharedMemTest << <grid, block, 4*1024>> >();
        cudaDeviceSynchronize();
    }

输出结果如下：

blockIdx.x is 27 s is at 1001000, ss is at 1000000 
blockIdx.x is 6 s is at 1001000, ss is at 1000000 
blockIdx.x is 9 s is at 1001000, ss is at 1000000 
blockIdx.x is 18 s is at 1001000, ss is at 1000000 
… 
blockIdx.x is 30 s is at 1001000, ss is at 1000000 
blockIdx.x is 10 s is at 1001000, ss is at 1000000

可以看出以下几点：
1，每个block 都有自己独立的shared memory地址空间。
2，静态开辟的空间总是从地址1000000开始。
3，动态开辟空间是在静态空间之后的。

如果将动态开辟地址大小设置太大，导致整个block 使用的shared memory 空间超过maxSharedMemoryPerBlock，会导致kernel 不执行。例如将调用代码改成下面：

{
        dim3 block(32);
        dim3 grid(32);
        sharedMemTest << <grid, block, 48*1024>> >();
        cudaDeviceSynchronize();
    }

由于我的显卡中maxSharedMemoryPerBlock = 48KB，动态空间+静态 = 49KB所以程序并没有输出。

本文来自朝气的CSDN 博客，全文地址请点击：https://blog.csdn.net/galaxy_wolf/article/details/50738200?utm_source=copy