GPU的矩陣轉置優化 transpose

2021-10-21 16:57:05 字數 2659 閱讀 5952

結論

在影象處理,深度學習領域,有很多矩陣運算的工作,而伴隨矩陣運算就存在大量的矩陣轉置,轉置不涉及計算,主要的工作都在資料的讀取寫入方面,所以如何加快資料搬移是一種很重要的優化點。

假設都是按行儲存。

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...