靜態分配

加上字首 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所以程式並沒有輸出。