CUDA ---- Shared Memory

時間 2019-11-14 標籤 cuda shared memory

CUDA SHARED MEMORY

shared memory在以前的博文有些介紹，這部分會專門講解其內容。在global Memory部分，數據對齊和連續是很重要的話題，當使用L1的時候，對齊問題能夠忽略，可是非連續的獲取內存依然會下降性能。依賴於算法本質，某些狀況下，非連續訪問是不可避免的。使用shared memory是另外一種提升性能的方式。程序員

GPU上的memory有兩種：算法

· On-board memory編程

· On-chip memory數組

global memory就是一塊很大的on-board memory，而且有很高的latency。而shared memory正好相反，是一塊很小，低延遲的on-chip memory，比global memory擁有高得多的帶寬。咱們能夠把他當作可編程的cache，其主要做用有：緩存

· An intra-block thread communication channel 線程間交流通道架構

· A program-managed cache for global memory data可編程cacheapp

· Scratch pad memory for transforming data to improve global memory access patterns函數

本文主要涉及兩個例子做解釋：reduction kernel，matrix transpose kernel。oop

shared memory（SMEM）是GPU的重要組成之一。物理上，每一個SM包含一個當前正在執行的block中全部thread共享的低延遲的內存池。SMEM使得同一個block中的thread可以相互合做，重用on-chip數據，而且可以顯著減小kernel須要的global memory帶寬。因爲APP能夠直接顯式的操做SMEM的內容，因此又被稱爲可編程緩存。性能

因爲shared memory和L1要比L2和global memory更接近SM，shared memory的延遲比global memory低20到30倍，帶寬大約高10倍。

當一個block開始執行時，GPU會分配其必定數量的shared memory，這個shared memory的地址空間會由block中的全部thread 共享。shared memory是劃分給SM中駐留的全部block的，也是GPU的稀缺資源。因此，使用越多的shared memory，可以並行的active就越少。

關於Program-Managed Cache：在C語言編程裏，循環（loop transformation）通常都使用cache來優化。在循環遍歷的時候使用從新排列的迭代順序能夠很好利用cache局部性。在算法層面上，咱們須要手動調節循環來達到使人滿意的空間局部性，同時還要考慮cache size。cache對於程序員來講是透明的，編譯器會處理全部的數據移動，咱們沒有能力控制cache的行爲。shared memory則是一個可編程可操做的cache，程序員能夠徹底控制其行爲。

Shared Memory Allocation

咱們能夠動態或者靜態的分配shared Memory，其聲明便可以在kernel內部也能夠做爲全局變量。

其標識符爲：__shared__。

下面這句話靜態的聲明瞭一個2D的浮點型數組：

__shared__ float tile[size_y][size_x];

若是在kernel中聲明的話，其做用域就是kernel內，不然是對全部kernel有效。若是shared Memory的大小在編譯器未知的話，可使用extern關鍵字修飾，例以下面聲明一個未知大小的1D數組：

extern __shared__ int tile[];

因爲其大小在編譯器未知，咱們須要在每一個kernel調用時，動態的分配其shared memory，也就是最開始說起的第三個參數：

kernel<<<grid, block, isize * sizeof(int)>>>(...)

應該注意到，只有1D數組才能這樣動態使用。

Shared Memory Banks and Access Mode

以前博文對latency和bandwidth有了充足的研究，而shared memory可以用來隱藏因爲latency和bandwidth對性能的影響。下面將解釋shared memory的組織方式，以便研究其對性能的影響。

Memory Banks

爲了得到高帶寬，shared Memory被分紅32（對應warp中的thread）個相等大小的內存塊，他們能夠被同時訪問。不一樣的CC版本，shared memory以不一樣的模式映射到不一樣的塊（稍後詳解）。若是warp訪問shared Memory，對於每一個bank只訪問很少於一個內存地址，那麼只須要一次內存傳輸就能夠了，不然須要屢次傳輸，所以會下降內存帶寬的使用。

Bank Conflict

當多個地址請求落在同一個bank中就會發生bank conflict，從而致使請求屢次執行。硬件會把這類請求分散到儘量多的沒有conflict的那些傳輸操做裏面，下降有效帶寬的因素是被分散到的傳輸操做個數。

warp有三種典型的獲取shared memory的模式：

· Parallel access：多個地址分散在多個bank。

· Serial access：多個地址落在同一個bank。

· Broadcast access：一個地址讀操做落在一個bank。

Parallel access是最一般的模式，這個模式通常暗示，一些（也多是所有）地址請求可以被一次傳輸解決。理想狀況是，獲取無conflict的shared memory的時，每一個地址都在落在不一樣的bank中。

Serial access是最壞的模式，若是warp中的32個thread都訪問了同一個bank中的不一樣位置，那就是32次單獨的請求，而不是同時訪問了。

Broadcast access也是隻執行一次傳輸，而後傳輸結果會廣播給全部發出請求的thread。這樣的話就會致使帶寬利用率低。

下圖是最優狀況的訪問圖示：

下圖一種隨機訪問，一樣沒有conflict：

下圖則是某些thread訪問到同一個bank的狀況，這種狀況有兩種行爲：

· Conflict-free broadcast access if threads access the same address within a bank

· Bank conflict access if threads access different addresses within a bank

Access Mode

根據不一樣的CC版本，bank的配置也不一樣，具體爲：

· 4 bytes for devices of CC 2.x

· 8 bytes for devices of CC3.x

對於Fermi，一個bank是4bytes。每一個bank的帶寬是32bits每兩個cycle。連續的32位字映射到連續的bank中，也就是說，bank的索引和shared memory地址的映射關係以下：

bank index = (byte address ÷ 4 bytes/bank) % 32 banks

下圖是Fermi的地址映射關係，注意到，bank中每一個地址相差32，相鄰的word分到不一樣的bank中以便使warp可以得到更多的並行獲取內存操做（獲取連續內存時，連續地址分配到了不一樣bank中）。

當同一個warp的兩個thread要獲取同一個地址（注意是同一個地址仍是同一個bank）的時候並不發生bank conflict。對於讀操做，會用一次transaction得到結果後廣播給全部請求，當寫操做時，只有一個thread會真正去寫，可是哪一個thread執行了寫是沒法知道的（undefined）。

在8bytes模式中，同理4bytes，連續的64-bits字會映射到連續的bank。每一個bank帶寬是64bite/1個clock。其映射關係公式：

bank index = (byte address ÷ 8 bytes/bank) % 32 banks

這裏，若是兩個thread訪問同一個64-bit中的任意一個兩個相鄰word（1byte）也不會致使bank conflict，由於一次64-bit（bank帶寬64bit/cycle）的讀就能夠知足請求了。也就是說，同等狀況下，64-bit模式通常比32-bit模式更少碰到bank conflict。

下圖是64-bit的關係圖。儘管word0和word32都在bank0中，同時讀這兩個word也不會致使bank conflict（64-bit/cycle）：

下圖是64-bit模式下，conflict-free的狀況，每一個thread獲取不一樣的bank：

下圖是另外一種conflict-free狀況，兩個thread或獲取同一個bank中的word：

下圖紅色箭頭是bank conflict發生的狀況：

Memory Padding

memory padding是一種避免bank conflict的方法，以下圖所示，全部的thread分別訪問了bank0的五個不一樣的word，這時就會致使bank conflict，咱們採起的方法就是在每N（bank數目）個word後面加一個word，這樣就以下面右圖那樣，本來bank0的每一個word轉移到了不一樣的bank中，從而避免了bank conflict。

增長的這寫word不會用來存儲數據，其惟一的做用就是移動原始bank中的word，使用memory padding會致使block可得到shared memory中有用的數量減小。還有就是，要從新計算數組索引來獲取正確的數據元素。

Access Mode Configuration

對Kepler來講，默認狀況是4-byte模式，能夠用下面的API來查看：

cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig *pConfig);

返回結果放在pConfig中，其結果能夠是下面兩種：

cudaSharedMemBankSizeFourByte

cudaSharedMemBankSizeEightByte

可使用下面的API來設置bank的大小：

cudaError_t cudaDeviceSetSharedMemConfig(cudaSharedMemConfig config);

bank的配置參數以下三種：

cudaSharedMemBankSizeDefault

cudaSharedMemBankSizeFourByte

cudaSharedMemBankSizeEightByte

在其啓動不一樣的kernel之間修改bank配置會有一個隱式的device同步。修改shared memory的bank大小不會增長shared memory的利用或者影響kernel的Occupancy，可是對性能是一個主要的影響因素。一個大的bank會產生較高的帶寬，可是鑑於不一樣的access pattern，可能致使更多的bank conflict。

Synchronization

由於shared Memory能夠被同一個block中的不一樣的thread同時訪問，當同一個地址的值被多個thread修改就致使了inter-thread conflict，因此咱們須要同步操做。CUDA提供了兩類block內部的同步操做，即：

· Barriers

· Memory fences

對於barrier，全部thread會等待其餘thread到達barrier point；對於Memory fence，全部thread會阻塞到全部修改Memory的操做對其餘thread可見，下面解釋下CUDA須要同步的主要緣由：weakly-ordered。

Weakly-Ordered Memory Model

現代內存架構有很是寬鬆的內存模式，也就是意味着，Memory的獲取沒必要按照程序中的順序來執行。CUDA採用了一種叫作weakly-ordered Memory model來獲取更激進的編譯器優化。

GPU thread寫數據到不一樣的Memory的順序（好比shared Memory，global Memory，page-locked host memory或者另外一個device上的Memory）一樣不必跟程序裏面順序呢相同。一個thread的讀操做的順序對其餘thread可見時也可能與實際上執行寫操做的thread順序不一致。

爲了顯式的強制程序以一個確切的順序運行，就須要用到fence和barrier。他們也是惟一能保證kernel對Memory有正確的行爲的操做。

Explicit Barrier

同步操做在咱們以前的文章中也提到過很多，好比下面這個：

void __syncthreads();

__syncthreads就是做爲一個barrier point起做用，block中的thread必須等待全部thread都到達這個point後才能繼續下一步。這也保證了全部在這個point以前獲取global Memory和shared Memory的操做對同一個block中全部thread可見。__syncthreads被用來協做同一個block中的thread。當一些thread獲取Memory相同的地址時，就會致使潛在的問題（讀後寫，寫後讀，寫後寫）從而引發未定義行爲狀態，此時就可使用__syncthreads來避免這種狀況。

使用__syncthreads要至關當心，只有在全部thread都會到達這個point時才能夠調用這個同步，顯而易見，若是同一個block中的某些thread永遠都到達該點，那麼程序將一直等下去，下面代碼就是一種錯誤的使用方式：

if (threadID % 2 == 0) {
    __syncthreads();
    } else {
        __syncthreads();
}

Memory Fence

這種方式保證了任何在fence以前的Memory寫操做對fence以後thread均可見，也就是，fence以前寫完了，fence以後其它thread就都知道這塊Memory寫後的值了。fence的設置範圍比較廣，分爲：block，grid和system。

能夠經過下面的API來設置fence：

void __threadfence_block();

看名字就知道，這個函數是對應的block範圍，也就是保證同一個block中thread在fence以前寫完的值對block中其它的thread可見，不一樣於barrier，該function不須要全部的thread都執行。

下面是grid範圍的API，做用同理block範圍，把上面的block換成grid就是了：

void __threadfence();

下面是system的，其範圍針對整個系統，包括device和host：

void __threadfence_system();

Volatile Oualifier

聲明一個使用global Memory或者shared Memory的變量，用volatile修飾符來修飾該變量的話，會組織編譯器作一個該變量的cache的優化，使用該修飾符後，編譯器就會認爲該變量可能在某一時刻被別的thread改變，若是使用cache優化的話，獲得的值就缺少時效，所以使用volatile強制每次都到global 或者shared Memory中去讀取其絕對有效值。

CHECKING THE DATA LAYOUT OF SHARED MEMORY

該部分會試驗一些使用shared Memory的例子，包括如下幾個方面：

· 方陣vs矩陣數組

· Row-major vs column-major access

· 靜態vs動態shared Memory聲明

· 全局vs局部shared Memory

· Memory padding vs no Memory padding

咱們在設計使用shared Memory的時候應該關注下面的信息：

· Mapping data elements across Memory banks

· Mapping from thread index to shared Memory offset

搞明白這兩點，就能夠掌握shared Memory的使用了，從而構建出牛逼的代碼。

Square Shared Memory

下圖展現了一個每一維度有32個元素並以row-major存儲在shared Memory，圖的最上方是該矩陣實際的一維存儲圖示，下方的邏輯的二維shared Memory：

咱們可使用下面的語句靜態聲明一個2D的shared Memory變量：

__shared__ int tile[N][N];

可使用下面的方式來數據，相鄰的thread獲取相鄰的word：

tile[threadIdx.y][threadIdx.x]

tile[threadIdx.x][threadIdx.y]

上面兩種方式哪一個更好呢？這就須要注意thread和bank的映射關係了，咱們最但願看到的是，同一個warp中的thread獲取的是不一樣的bank。同一個warp中的thread可使用連續的threadIdx.x來肯定。不一樣bank中的元素一樣是連續存儲的，以word大小做爲偏移。所以次，最好是讓連續的thread（由連續的threadIdx.x肯定）獲取shared Memory中連續的地址，由此得知，

tile[threadIdx.y][threadIdx.x]應該展示出更好的性能以及更少的bank conflict。

Accessing Row-Major versus Column-Major

假設咱們的grid有2D的block（32,32），定義以下：

#define BDIMX 32
#define BDIMY 32
dim3 block(BDIMX,BDIMY);
dim3 grid(1,1);

咱們對這個kernel有以下兩個操做：

· 將thread索引以row-major寫到2D的shared Memory數組中。

· 從shared Memory中讀取這些值並寫入到global Memory中。

kernel代碼：

__global__ void setRowReadRow(int *out) {
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX];
    // 由於block只有一個
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;
    // 這裏同步是爲了使下面shared Memory的獲取以row-major執行
    //如有的線程未完成，而其餘線程已經在讀shared Memory。。。
    __syncthreads();
    // shared memory load operation
    out[idx] = tile[threadIdx.y][threadIdx.x] ;
}

觀察代碼可知，咱們有三個內存操做：

· 向shared Memory存數據

· 從shared Memor取數據

· 向global Memory存數據

由於在同一個warp中的thread使用連續的threadIdx.x來檢索title，該kernel是沒有bank conflict的。若是交換上述代碼threadIdx.y和threadIdx.x的位置，就變成了column-major的順序。每一個shared Memory的讀寫都會致使Fermi上32-way的bank conflict或者在Kepler上16-way的bank conflict。

__global__ void setColReadCol(int *out) {
    // static shared memor
    __shared__ int tile[BDIMX][BDIMY];
    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    // shared memory store operation
    tile[threadIdx.x][threadIdx.y] = idx;
    // wait for all threads to complete
    __syncthreads();
    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

編譯運行：

$ nvcc checkSmemSquare.cu –o smemSquare
$ nvprof ./smemSquare

在Tesla K40c（4-byte模式）上的結果以下，正如咱們所想的，row-major表現要出色：

./smemSquare at device 0 of Tesla K40c with Bank Mode:4-byte
<<< grid (1,1) block (32,32)>>
Time(%) Time Calls Avg Min Max Name
13.25% 2.6880us 1 2.6880us 2.6880us 2.6880us setColReadCol(int*)
11.36% 2.3040us 1 2.3040us 2.3040us 2.3040us setRowReadRow(int*)

而後使用nvprof的下面的兩個參數來衡量相應的bank-conflict：

shared_load_transactions_per_request

shared_store_transactions_per_request

結果以下（8 bytes模式，4 bytes應該是32），row-major只有一次transaction：

Kernel:setColReadCol (int*)
1 shared_load_transactions_per_request 16.000000
1 shared_store_transactions_per_request 16.000000
Kernel:setRowReadRow(int*)
1 shared_load_transactions_per_request 1.000000
1 shared_store_transactions_per_request 1.000000
Writing Row-Major and Reading Column-Major

本節的kernel實現以row-major寫shared Memory，以Column-major讀shared Memory，下圖指明瞭這兩種操做的實現：

kernel代碼：

__global__ void setRowReadCol(int *out) {
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX];
    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;
    // wait for all threads to complete
    __syncthreads();
    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

查看nvprof結果：

Kernel:setRowReadCol (int*)
1 shared_load_transactions_per_request 16.000000
1 shared_store_transactions_per_request 1.000000

寫操做是沒有conflict的，讀操做則引發了一個16次的transaction。

Dynamic Shared Memory

正如前文所說，咱們能夠全局範圍的動態聲明shared Memory，也能夠在kernel內部動態聲明一個局部範圍的shared Memory。注意，動態聲明必須是未肯定大小一維數組，所以，咱們就須要從新計算索引。由於咱們將要以row-major寫，以colu-major讀，因此就須要保持下面兩個索引值：

· row_idx：1D row-major 內存的偏移

· col_idx：1D column-major內存偏移

kernel代碼：

__global__ void setRowReadColDyn(int *out) {
    // dynamic shared memory
    extern __shared__ int tile[];
    // mapping from thread index to global memory index
    unsigned int row_idx = threadIdx.y * blockDim.x + threadIdx.x;
    unsigned int col_idx = threadIdx.x * blockDim.y + threadIdx.y;
    // shared memory store operation
    tile[row_idx] = row_idx;
    // wait for all threads to complete
    __syncthreads();
    // shared memory load operation
    out[row_idx] = tile[col_idx];
}

kernel調用時配置的shared Memory：

setRowReadColDyn<<<grid, block, BDIMX * BDIMY * sizeof(int)>>>(d_C);

查看transaction：

Kernel: setRowReadColDyn(int*)
1 shared_load_transactions_per_request 16.000000
1 shared_store_transactions_per_request 1.000000

該結果和以前的例子相同，不過這裏使用的是動態聲明。

Padding Statically Declared Shared Memory

直接看kernel代碼：

__global__ void setRowReadColPad(int *out) {
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX+IPAD];
    // mapping from thread index to global memory offset
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;
    // wait for all threads to complete
    __syncthreads();
    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

改代碼是setRowReadCol的翻版，查看結果：

Kernel: setRowReadColPad(int*)
1 shared_load_transactions_per_request 1.000000
1 shared_store_transactions_per_request 1.000000

正如指望的那樣，load的bank_conflict已經消失。在Fermi上，只須要加上一列就能夠解決bank-conflict，可是在Kepler上卻不必定，這取決於2D shared Memory的大小，所以對於8-byte模式，可能須要屢次試驗才能獲得正確結果。

參考書《professional cuda c programming》