結論
在影象處理,深度學習領域,有很多矩陣運算的工作,而伴隨矩陣運算就存在大量的矩陣轉置,轉置不涉及計算,主要的工作都在資料的讀取寫入方面,所以如何加快資料搬移是一種很重要的優化點。假設都是按行儲存。
int row =
1024
int col =
512void
transpose_cpu
(vector<
int>
& in, vector<
int>
& out)
}}
//版本1
//假如是block = row; grid = col
__global__
void
transpose_cuda_kernel
(int
* in,
int* out)
//版本2
//假如grid和block的設定都為二維(x,y),矩陣大小為(nrows,ncols)
首先來看一下上面的圖,這個綠的block中有(3,4)個執行緒,在gpu中block的thread是以warp為單位執行的,每個warp為32個執行緒,這32執行緒執行同樣的指令,如果說這32個執行緒訪問的記憶體的連續,那麼就可以合併為乙個資料處理事務(具體看資料大小),首先執行緒去l1/l2快取中去查詢,如果不存在就去全域性記憶體中取數放到快取中。如果說訪問的全域性記憶體位址不是連續的,那麼資料處理事務就會分為多次處理,所以最好能形成資料合併,在讀取中可以看到,執行緒訪問的記憶體是連續,但是在寫資料的時候是按列去寫的,但是記憶體是按行布局的,所以資料訪存無法合併,那麼如何解決呢?
我們的目的就是讓記憶體能夠連續寫入,這裡引入share memory, 這裡邏輯有點繞,為了搞清楚需要好好捋一下,首先看看我們的目的,我們是為了執行緒在寫入記憶體的時候,連續的執行緒能夠訪問連續的記憶體位址,如下圖中的set num 的效果
先看看原始是什麼情況, 由於記憶體是按行儲存的,所以在set num的時候,連續的執行緒拿到的輸出位址不是連續的,這樣本來三個執行緒在取數的時候可以合併乙個指令取出,結果在賦值的時候三個執行緒卻不能合併指令,因為三個位址不連續,示意圖如下:
所以如何解決在賦值的時候連續執行緒可以對應連續位址就好了,在計算機中有個經典的思路,所有的問題都可以通過乙個中介軟體去解決問題,在這裡就是利用share memory.
**
第一步:
//假設矩陣大小為(m,n)
int in_x = blockidx.x * blockdim.x + threadidx.x;
int in_y = blockidx.y * blockdim.y + threadidx.y;
tmp[threadidx.y]
[threadidx.x]
= in[in_y * n + in_x]
;// 這裡可以看到當threadidx連續時候in[pos]中pos也連續
第二步:
int bid = threadidx.y*blockdim.x + threadidx.x;
// block拉成一維後
x = bid/blockdim.y;
//對應的share memory的位置
y = bid%blockdim.y;
int to_x = blockidx.y * blockdim.y + y;
// blockidx.y * blockdim.y這是block的新位置
int to_y = blockidx.x * blockdim.x + x;
// transpose with boundary test
out[to_y * m + to_x]
= tmp[y]
[x];
// 這裡可以看到當threadidx連續的時候out[pos]中pos也連續
總的**如下
#define index(row, col, inter) ((row) * (inter) + (col))
__global__
void *****_cuda_kernel(float *out, float *in, const int nrows, const int ncols)
}
從實驗結果來看,頻寬利用率提公升明 稀疏矩陣的轉置(矩陣轉置和快速轉置)
實現矩陣的轉置 1.將矩陣的行列值相互轉換。2.將每個三元組中的i和j交換。3.重排三元組之間的次序便可實現矩陣的轉置。void transposesmatrix tsmatrix m,tsmatrix t return transposesmatrix快速轉置的原理是 如果能預先確定矩陣m中每一列...
稀疏矩陣轉置矩陣
num 矩陣a中某列的非零元素的個數 cpot 初值表示矩陣a中某列第乙個非零元素在b中的位置,並有如下遞推 cpot 1 0 cpot col cpot col 1 num col 1 2 col nu end right.const int maxterm 100 struct sparsema...
稀疏矩陣的轉置與快速轉置
假設在m n的矩陣中,有t個元素不為0。令稀疏因子s t m n 通常認為s 0.05時稱為稀疏矩陣。有時為了節省儲存空間,可以對這類矩陣進行壓縮儲存。所謂的壓縮儲存就是,為多個相同的值分配儲存在乙個空間,對零元不分配空間。而稀疏矩陣是只儲存有效值,無效值只分配乙個空間。在這裡我們用乙個順序表vec...