CUDA之靜態 動態共享記憶體分配詳解

2021-07-29 13:00:40 字數 2502 閱讀 7611

#include __global__ void staticreverse(int *d, int n)


__global__ void dynamicreverse(int *d, int n)


int main(void)


int *d_d;


cudamalloc(&d_d, n * sizeof(int));


// run version with static shared memory


cudamemcpy(d_d, a, n*sizeof(int), cudamemcpyhosttodevice);


staticreverse<<<1,n>>>(d_d, n);


cudamemcpy(d, d_d, n*sizeof(int), cudamemcpydevicetohost);


for (int i = 0; i < n; i++)


if (d[i] != r[i]) printf("error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);


// run dynamic shared memory version


cudamemcpy(d_d, a, n*sizeof(int), cudamemcpyhosttodevice);


dynamicreverse<<<1,n,n*sizeof(int)>>>(d_d, n);


cudamemcpy(d, d_d, n * sizeof(int), cudamemcpydevicetohost);


for (int i = 0; i < n; i++)


if (d[i] != r[i]) printf("error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);


}
上面的**使用共享儲存器對大小為64的陣列進行逆序處理。這兩個核函式十分相似，不同之處在於共享記憶體陣列的宣告以及核函式的呼叫。

如果共享記憶體陣列的大小在編譯時就可以確定，就像在上節**中staticreverse核函式中寫的那樣，我們就可以顯式地宣告固定大小的陣列，下面是我們宣告的s陣列：

__global__ void staticreverse(int *d, int n)
在這個核函式中，t和tr分別代表了原始和倒序之後陣列的下標索引。每個執行緒使用語句s[t] = d[t]將全域性記憶體的資料拷貝到共享記憶體，反向工作是通過語句d[t] = s[tr]來完成的。但是在執行執行緒訪問共享記憶體中被執行緒寫入的資料前，記住要使用__syncthreads()來確保所有的執行緒都已經完全將資料載入到共享記憶體。

在這個例子中，使用共享記憶體是用於促進全域性記憶體合併訪問(在舊的cuda裝置上，計算能力1.1或更低)。對於讀取和寫入都實現了最優的全域性儲存器合併，因為全域性記憶體總是通過線性對齊的索引t來訪問的。反向索引tr僅用於訪問共享儲存器，其不具有全域性儲存器的順序訪問限制，因此不能獲得最佳效能。共享記憶體的唯一效能問題是bank衝突，我們之後會做討論。

另乙個核函式使用了動態分配共享記憶體的方式，這主要用於共享記憶體的大小在編譯時不能確定的情況。在這種情況下，每個執行緒塊中共享記憶體的大小必須在核函式第三個執行配置引數中指定(以位元組為單位)，如下所示：

dynamicreverse<<<1, n, n*sizeof(int)>>>(d_d, n);
該動態共享記憶體的核函式dynamicreverse()使用了未指定大小的extern陣列語法(extern __shared__ int s)來宣告共享記憶體陣列。

note:注意中括號與extern說明符。

當核函式被啟動時，陣列大小從第三個執行配置引數被隱式地確定。該核函式其餘部分的**與staticreverse()核函式相同。

而如果你想在乙個核函式中動態地申請多個陣列時該怎麼辦呢？你必須在首先申請乙個單獨的未指定大小的extern陣列，然後使用指標將它分為多個陣列，如下所示：

extern __shared__ int s;


int *integerdata = s;                        // ni ints


float *floatdata = (float*)&integerdata[ni]; // nf floats


char *chardata = (char*)&floatdata[nf];      // nc chars
這樣的話，你需要在核函式中這樣指定共享記憶體的大小：

mykernel<<>>(...);

CUDA申請動態共享記憶體

直接上 include include global void kernel int main return 0 程式輸出 1021.000000 1022.000000 1023.000000 index 1021.000000 index 1022.000000 index 1023.00000...

靜態記憶體分配和 動態記憶體分配

1 靜態記憶體分配是在編譯時完成的，不需要占用cpu資源 動態分配記憶體是在執行時完成的，動態記憶體的分配與釋放需要占用cpu資源 2 靜態記憶體分配是在棧上分配的，動態記憶體是堆上分配的 3 動態記憶體分配需要指標或引用資料型別的支援，而靜態記憶體分配不需要 4 靜態分配記憶體需要在編譯前確定記憶...

靜態分配記憶體和動態分配記憶體

首先，在使用動態分配記憶體技術前，必須明白自己在做什麼，這樣做與其它的方法有什麼不同，特別是會產生哪些負面影響，天下沒有免費的午餐。動態分配記憶體與靜態分配記憶體的區別 1 靜態記憶體分配是在編譯時完成的，不需要占用cpu資源 動態分配記憶體是在執行時完成的，動態記憶體的分配與釋放需要占用cpu資源...