CUDA（三）三種memory的活用

分類：CUDA

http://blog.csdn.net/wesley_2013/article/details/11948761

（465）（0）舉報收藏

此部落格為博主的自學筆記
，歡迎大家共同交流，如果有錯誤的地方歡迎留言指正。

GPU程式設計時常用的高速訪問記憶體有三種，分別為：register, shared memory, constant memory。

以下面一個程式段為例：

設： A_d[ 256 ] , B_d[ 256 * 256 ], C_d[256]為global memory的記憶體空間。

kernel函式中的程式語句：

for(int i = 0; i < 256; i ){

A_d[ threadIdx.x ] = A_d [ threadIdx.x ] B_d[ threadIdx.x * 256 i ] * C_d[ i ];

}

這個是將B,C矩陣相乘再將結果存到A中。

REGISTER的活用

很明顯，上面的寫法對於A距陣的每一個元素進行計算時，都要進行256次global memory記憶體空間的訪問。而global memory的訪問速度又非常的緩慢，大大降低了效能。

但是，通過register的活用，我們可以輕鬆解決這個問題。因為上面的例子，執行緒之間沒有交叉的記憶體訪問，所以，我們可以將它的迴圈中的中間結果儲存於register中的變數中。這樣可以大大地加快訪問速度，從而提高效能。

float temp = 0;

for(int i = 0; i < 256; i ){

temp = temp
B_d[ threadIdx.x * 256 i ] * C_d[ i ];

}

A_d[ threadIdx.x ] = temp;

SHARED MEMORY的活用

不難看出，上面的程式語句中，C_d只有一維陣列，所佔記憶體空間較小，且被頻繁訪問，這樣我們可以將他複製到空間有限的shared memory當中，加快訪問的速度，以提高效能。

__shared__ float C[256];

C[ threadIdx.x] = C_d[ threadIdx.x];

for(int i = 0; i < 256; i ){

A_d[ threadIdx.x ] = A_d [ threadIdx.x ] B_d[ threadIdx.x * 256 i ] * C[
i ];

}

CONSTANT MEMORY的活用

如果shared memory的空間不足，我們可以使用其它訪問速度較快的記憶體來代替。因為此次的C_d只是被當做乘數來使用，無需要修改，所以我們可以將他儲存在同樣訪問速度較快的只讀constant memory中。只是他的宣告和初始化同shared memory有些不同，不是在kernel函式中，而是主機側。

主機側：

__constant__ float C_c[ 256 ];

初始化陣列C之後，通過以下語句將C複製到C_c中

cudaMemcpyToSymbol(C_c, C, 256*sizeof(float));

GPU側：

for(int i = 0; i < 256; i ){

A_d[ threadIdx.x ] = A_d [ threadIdx.x ] B_d[ threadIdx.x * 256 i ] * C_c[
i ];

}

通過上述步驟，計算效能會有很大的提高。

儘量充分利用這三種記憶體訪問方式，以便發揮出GPU的最大效能。

CUDA（三） 三種memory的活用