2.2CUDA-Memory(存儲)和bank-conflict

在CUDA基本概念介紹有簡單介紹CUDA memory。這里詳細介紹：

每一個線程擁有自己的私有存儲器，每一個線程塊擁有一塊共享存儲器(Shared memory)；最后，grid中所有的線程都可以訪問同一塊全局存儲器(global memory)。除此之外，還有兩種可以被所有線程訪問的只讀存儲器：常數存儲器（constant memory）和紋理存儲器(Texture memory),它們分別為不同的應用進行了優化。全局存儲器、常數存儲器和紋理存儲器中的值在一個內核函數執行完成后將被繼續保持，可以被同一程序中其也內核函數調用。

下表給出了這8種存儲器的位置、緩存情況，訪問權限及生存域

 

存儲器	位置	擁有緩存	訪問權限	變量生存周期
register	GPU片內	N/A	Device可讀/寫	與thread相同
Local memory	板載顯存	無	Device可讀/寫	與thread相同
Shared memory	GPU片內	N/A	Device可讀/寫	與block相同
Constant memory	板載顯存	有	Device可讀，host可讀寫	可在程序中保持
Texture memory	板載顯存	有	Device可讀，host可讀寫	可在程序中保持
Global memory	板載顯存	無	Device可讀/寫, host可讀/寫	可在程序中保持
Host memory	Host內存	無	host可讀/寫	可在程序中保持
Pinned memory	Host內存	無	host可讀/寫	可在程序中保持

 

kernel變量定義，使用范圍和生命周期。

其中__shared__和__constant__前面的__device__聲明是可以省略的

 

Global memory， 如果有一個thread修改啦global memory的值，其他的thread不能立即看到這個值的變化。需要終止這個kernel，然后lanuch一個新的kernel，這樣新的kernel能看到global memory的變化。

 

Shared Memory（共享存儲）

由於訪問速度比Global快的多，比如向量加法，每次從in指針里面取內容都是從global里面取。比如矩陣乘法：

__global__ void MatrixMulKernel(int m, int n, int k, float* A, float*
B, float* C)
{
    int Row = blockIdx.y*blockDim.y+threadIdx.y;
    int Col= blockIdx.x*blockDim.x+threadIdx.x;
    if ((Row < m) && (Col < k)) {
        float Cvalue = 0.0;
        for (int i = 0; i < n; ++i)
            /* A[Row, i] and B[i, Col] */
            Cvalue += A[Row*n+i] * B[Col+i*k];
        C[Row*k+Col] = Cvalue;
    }
}
                        

可以看到有很多值取拉多次,重復多次訪問global memory。我們可以把需要用的數據保存在shared memory中，如下圖：

減少啦訪問的次數，而且訪問shared memory 速度更快

總結一個公共的編程策略模型：

當我們在划分存儲數據時經常划分成很多塊，或者是tile，在這里稱作 partition data or tile data.

1. 把partition data or tile data 保存在 shared memory里面。

2. 執行計算時，從shared memory里面取這些數據。

3. 上面的結束之后，拷貝shared memory的數據到global memory

 

合並訪問

warp是調度和執行的基本單位，這個在上一篇中有提到，half-warp是存儲器操作的基本單位，這兩個非常重要。

到我們都知道每一個half-warp是16個thread.以tesla為例：

Tesla 的每個 SM 擁有 16KB 共享存儲器,用於同一個線程塊內的線程間通信。為了使一個 half-warp 內的線程能夠在一個內核周期中並行訪問,共享存儲器被組織成 16 個 bank,

每個 bank 擁有1024Kb 的寬度,一個Int 4個Byte ,故每個 bank 可保存 256 個整形或單精度浮點數,或者說目前的bank 組織成了 256 行 16 列的矩陣

 

 上圖中shared memory的長度是256.

舉例說明：

__shared__ int data[128];

那么data[0], data[1]...data[15] 會依次訪問bank[0],bank[1]...bank[15].

而data[16] ...data[31] 又會以此訪問bank0 ...bank15. 

 由於存取內存是half-warp=16，所以屬於不同half-warp的thread不存在bank conflict.

因此,如果程序在存取 shared memory 的時候,使用以下的方式:

int number = data[base + tid]; (這個是連續訪問的，和base沒什么關系)

那就不會有任何 bank conflict,可以達到最高的效率。但是,如果是以下的方式:

int number = data[base + 4 * tid];

那么,thread 0 和 thread 4 就會存取到同一個 bank,thread1 和 thread 5 也是同 樣,這樣就會造成 bank conflict。在這個例子中,一個 half warp 的 16 個 threads 會有四個threads 存取同一個 bank,因此存取 share memory 的速度會變成原來的 1/4。

下面這種情況比較特殊：

int number = data[3]. 

大家都訪問同一個bank的同一個數據的時候，就可以形成一個broadcast，那樣就會把數據同時廣播給16個thread，這樣就可以合理利用shared memory的broadcast的機制。

 

解決bank conflict的策略

很多時候 shared memory 的 bank conflict 可以透過修改數據存放的方式來解決。例如,以下的程序:

data[tid] = global_data[tid];
...
int number = data[16 * tid]; 
會造成嚴重的 bank conflict,為了避免這個問題,可以把數據的排列方式稍加修改,把存取方式改成:

int row = tid / 16;
int column = tid % 16;
data[row * 17 + column] = global_data[tid];
...
int number = data[17 * tid];
這樣就不會造成 bank conflict 了。

簡單的說,矩陣中的數據是按照bank存儲的,第i個數據存儲在第i%16個bank中。一個block要訪問shared memory,只要能夠保證以其中相鄰的16個線程一組訪問thread,每個線程與bank是一一對應就不會產生bank conflict。否則會產生bankconflict,訪存時間成倍增加,增加的倍數由一個bank最多被多少個thread同時訪問決定。有一種極端情況,就是所有的16個thread同時訪問同一bank時反而只需要一個訪問周期,此時產生了一次廣播。

 

下面有一些小技巧可以避免bank conflict 或者提高global存儲器的訪問速度

1. 盡量按行操作,需要按列操作時可以先對矩陣進行轉置

2. 划分子問題時,使每個block處理的問題寬度恰好為16的整數倍,使得訪存可以按照 s_data[tid]=i_data[tid]的形式進行

3. 使用對齊的數據格式,盡量使用nvidia定義的格式如float3,int2等,這些格式本身已經對齊。

4. 當要處理的矩陣寬度不是16的整數倍時,將其補為16的整數倍,或者用malloctopitch而不是malloc。

5. 利用廣播,例如s_odata[tid] = tid%16 < 8 ? s_idata[tid] :s_idata[15];會產生8路的塊訪問沖突而用:

s_odata[tid]=s_idata[15];s_odata[tid]= tid%16 < 8 ? s_idata[tid] :s_data[tid]; 則不會產生塊訪問沖突