cuda shared memory 靜態分配和動態分配
阿新 • • 發佈:2018-12-12
靜態分配
加上字首 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所以程式並沒有輸出。