NO IMAGE

CUDA(三) 三種memory的活用

   

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

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的最大效能。