2. 程式人生 > >CUDA 共享記憶體 bank conflict

CUDA 共享記憶體 bank conflict

阿新 • • 發佈：2018-11-11

1. bank conflict

本文所有的實驗針對 GTX980 顯示卡，Maxwell 架構，計算能力 5.2。

GPU 共享記憶體是基於儲存體切換的架構（bank-switched-architecture）。在 Femi，Kepler，Maxwell 架構的裝置上有 32 個儲存體（也就是常說的共享記憶體分成 32 個bank），而在 G200 與 G80 的硬體上只有 16 個儲存體。

每個儲存體（bank）每個週期只能指向一次操作（一個 32bit 的整數或者一個單精度的浮點型資料），一次讀或者一次寫，也就是說每個儲存體（bank）的頻寬為 每週期 32bit。

如下圖所示，在一個執行緒塊中申請如下的共享記憶體：

__shared__ float sData[32][32];
  • 1

這裡寫圖片描述

也就是說在上述的 32 * 32 的二維陣列共享記憶體中，每一列對應同一個 bank。

  1. 同常量記憶體一樣，當一個 warp 中的所有執行緒訪問同一地址的共享記憶體時，會觸發一個廣播（broadcast）機制到 warp 中所有執行緒，這是最高效的。
  2. 如果同一個 warp 中的執行緒訪問同一個 bank 中的不同地址時將發生 bank conflict。
  3. 每個 bank 除了能廣播（broadcast）還可以多播（mutilcast）（計算能力 >= 2.0），也就是說，如果一個 warp 中的多個執行緒訪問同一個 bank 的同一個地址時（其他執行緒也沒有訪問同一個bank 的不同地址）不會發生 bank conflict。
  4. 即使同一個 warp 中的執行緒 隨機的訪問不同的 bank，只要沒有訪問同一個 bank 的不同地址就不會發生 bank conflict。

這裡寫圖片描述

如上圖所示，左側和右側的都沒有發生 bank conflict。而中間的存在 bank conflcit，由於經過最多兩次，該 warp 中的執行緒就都可以得到所要的資料，所有稱為 2-way bank conflict，如果同一個 warp 中的所有執行緒訪問一個 bank 中的 32 個不同地址，則需要分 32 次，稱為 32-way bank conflict。

這裡寫圖片描述

如上圖所示，左中右均未發生 bank conflict。

依次我們可以總結：只要同一個 warp 的不同執行緒會訪問到同一個 bank 的不同地址就會發生 bank conflict，除此之外的都不會發生 bank conflict。

既然廣播是針對同一個 warp 而言的，那麼如果不同的 warp 訪問同一個 bank 中的同一個地址呢？由於 每個 SM 中有 4 個 warp scheduler （GTX980），可以很好的排程 warp，使其 warp 之間的訪問衝突可以充分的隱藏，因此對效率的影響很小，遠遠小於 warp 內的 bank conflict。至於 warp scheduler 的排程機制，NVIDIA 沒有說的特別清楚，可能也是想要開發者不要過於關注於此。

2. 實驗 1

實現定義如下圖所示的 32 * 32 執行緒塊，共 1024 個執行緒，32 個 warp。

這裡寫圖片描述

申請如 1 中所示的 32 * 32 的共享記憶體，共 32 個 bank，每個 bank 對應 32 個元素。

  • 實驗 1.1
    該執行緒塊中的每個 warp 讀寫不同的 bank，不同的 warp 不會訪問一個地址，也就是一一對應的關係。圖中的數字就表示上圖中的執行緒標號。經分析可知，此時是沒有 bank conflict 的。

這裡寫圖片描述

程式碼如下：

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列座標
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行座標
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.y][threadIdx.x];
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 實驗 1.2
    該執行緒塊中的每個 warp 讀寫相同的 bank 的不同地址，不同的 warp 訪問不同，也就是一一對應的關係。圖中的數字就表示上圖中的執行緒標號。經分析可知，此時是存在很嚴重的 bank conflict 。

這裡寫圖片描述

程式碼如下：

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列座標
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行座標
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if (x_id < col && y_id < row)
{
    sData[threadIdx.x][threadIdx.y] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.x][threadIdx.y];
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 實驗 1.3（避免 bank conflict 的技巧）
    針對實驗 1.2 中出現的嚴重的 bank conflict，我們可以通過新增一個附加列來避免 bank conflict，如下圖所示，左圖為申請的共享記憶體矩陣形式，右圖是表示成 bank 後的形式，通過這種方式，原來在一個 bank 中的同一個 warp 都正好偏移到了不同的 bank 中。

這裡寫圖片描述

程式碼如下：

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列座標
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行座標
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE+1];

if (x_id < col && y_id < row)
{
    sData[threadIdx.x][threadIdx.y] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.x][threadIdx.y];
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12

上述三個小實驗的執行時間為：

實驗 1.1 ：0.052416 ms
實驗 1.2 ：0.131072 ms
實驗 1.3 ：0.053280 ms
  • 1
  • 2
  • 3

除去公共程式碼後的時間為：

實驗 1.1 ：0.034816 ms
實驗 1.2 ：0.113472 ms
實驗 1.3 ：0.035680 ms
  • 1
  • 2
  • 3

結論：

  1. 通過額外的一行，可以避免 bank conflict，執行時間與完全沒有 bank conflict 的執行時間差距很小。
  2. 存在 bank conflict 的，執行時間幾乎是沒有 bank conflict 的執行時間的 4 倍。

其實只要新增的是奇數列就可以，只不過 1 列是最節省空間（共享記憶體太寶貴）的。

3. 實驗 2

  • 實驗 2.1
    同一個 block 中所有第 i 列的執行緒都計算第 i 行的元素的和，此時所有同一個warp 會訪問同一個 bank 的不同地址。如下圖所示，分別表示第 0 列訪問 bank 0 中的第一個地址，第 1 列訪問 bank 1 中的第 1 個地址，依次類推。

這裡寫圖片描述

程式碼如下：

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列座標
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行座標
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();

    float data = 0.0f;
    for (int j = 0; j < BLOCKSIZE; j++)
    {
        data += sData[threadIdx.x][j];
    }
    matrixTest[index] = data;
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 實驗 2.2
    同實驗 1.3 類似，新增額外的一列，如下圖所示：

這裡寫圖片描述

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列座標
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行座標
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE+1];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();

    float data = 0.0f;

    for (int j = 0; j < BLOCKSIZE; j++)
    {
        data += sData[threadIdx.x][j];
    }
    matrixTest[index] = data;
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 19

上述兩個實驗的執行時間如下所示：

實驗 2.1 ：0.458144 ms
實驗 2.2 ：0.090848 ms
  • 1
  • 2

從上圖也可以看出，修改後的頻寬相當於修改前的 32 倍。修改後的執行時間也明顯得到改善。

4. 實驗 3

  • 實驗 3.1
    採用實驗 1.1 的方式，同一個 warp 訪問不同的 bank，不同的 warp 訪問不同的地址。

程式碼如下：

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列座標
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行座標
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();

    float data = 0.0f;
    for (int j = 0; j < 1000; j++)
    {
        data = sData[threadIdx.y][threadIdx.x];
    }
    matrixTest[index] = data;
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 實驗 3.2
    同一個 warp 訪問不同的 bank，所有 warp 訪問同一個地址，也就是說所有的行都會訪問第 0 行。

程式碼如下：

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列座標
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行座標
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();

    float data = 0.0f;
    for (int j = 0; j < 1000; j++)
    {
        data = sData[0][threadIdx.x];
    }
    matrixTest[index] = data;
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18

上述兩個實驗的執行時間如下所示：

實驗 2.1 ：0.053800 ms
實驗 2.2 ：0.055328 ms
  • 1
  • 2

在實驗 2.2 中存在明顯的不同 warp 間的衝突，但是執行時間差距很小，也就是說 warp 間衝突的影響很小。

5. visual profiler

通過 visual profiler 可以判斷程式中是否存在 bank conflict，在執行 visual profiler 前需要新增 -lineinfo選項，在 visual studio 中可以設定，如下所示：

這裡寫圖片描述

在 visual profiler 中分析實驗 1.2，結果如下所示，可以直接定位到出現 bank conflict 的行。

這裡寫圖片描述

6. 完整程式碼

我的GitHub

7. 參考

  1. 《CUDA_C_Programming_Guide》7.0 Appendix G. COMPUTE CAPABILITIES / 4.1 / 4.2
  2. 《CUDA_C_Best_Practices_Guide》7.0
  3. 《CUDA 並行程式設計：GPU程式設計指南》6.4
  4. 《GPU 高效能運算之 CUDA》4.7.1.3/4.4.3
  5. 《Performance modeling of atomic additions on
    GPU scratchpad memory》
  6. 《stackoverflow》

--------------------- 作者：木子超同學 來源：CSDN 原文：https://blog.csdn.net/endlch/article/details/47043069?utm_source=copy 版權宣告：本文為博主原創文章，轉載請附上博文連結！

相關推薦

CUDA 共享記憶體 bank conflict

1. bank conflict 本文所有的實驗針對 GTX980 顯示卡，Maxwell 架構，計算能力 5.2。 GPU 共享記憶體是基於儲存體切換的架構（bank-switched-architecture）。在 Femi，Kepler，Maxwell 架構的裝置上有 3

CUDA共享記憶體的使用示例

1 #include <cuda.h> 2 #include <cuda_runtime.h> 3 #include <device_launch_parameters.h> 4 #include <device_functions.h> 5 #in

cuda共享記憶體，全域性記憶體，紋理等的解釋

      開始閱讀粗大資料感覺文獻1的描述講得明白，摘錄到這裡方便他人瞭解。 增加簡單排版後，摘錄 1.共享記憶體 目前 CUDA 裝置中，每個 multiprocessor 有 16KB 的 shared memory。 Shared memory 分成16 個 ba

CUDA程式設計（七）共享記憶體與Thread的同步

https://blog.csdn.net/sunmc1204953974/article/details/51078818   CUDA程式設計（七） 共享記憶體與Thread的同步 在之前我們通過block，繼續增大了執行緒的數量，結果還是比較令人滿意的，但是也產生了一

CUDA GPU程式設計如何避免Bank conflict

 CUDA GPU程式設計如何避免Bank conflict     強烈推薦參考書：Professional CUDA C Programming，本書第234頁開始對bank conflict有非常詳盡的講解，本文正是從本書的學習和自己的程式設計經驗中總

CUDA共享內存的使用示例

blocks col all square 歸約 如果 ont 位置 nload CUDA共享內存使用示例如下：參考教材《GPU高性能編程CUDA實戰》。P54-P65 教材下載地址：http://download.csdn.net/download/yizhaoyanbo

以佇列的形式使用共享記憶體

共享記憶體允許多個程序使用某一段儲存區，資料不需要在程序之間複製，多個程序都把該記憶體區域對映到自己的虛擬地址空間，直接訪問該共享記憶體區域，從而可以通過該區域進行通訊，是一種速度很快IPC。 下面是共享記憶體對映圖 一般描述使用共享記憶體，都是以普通緩衝區的形式訪問，這裡

windows 驅動開發 MDL 核心層 使用者層共享記憶體

參考資料    https://blog.csdn.net/wdykanq/article/details/7752909    http://blog.51cto.com/laokaddk/404584   核心層建立記憶體對映使用者層 PVOID

ython實現程序間的通訊有Queue,Pipe,Value+Array等,其中Queue實現多個程序間的通訊,而Pipe實現兩個程序間通訊,而Value+Array使用得是共享記憶體對映檔案的方式,所以速度比較快

1.Queue的使用 from multiprocessing import Queue,Process import os,time,random #新增資料函式 def proc_write(queue,urls): print("程序(%s)正在寫入..."%(os.getpid()))

Linux-程序通訊-訊息佇列/訊號燈/共享記憶體

訊息佇列     訊息佇列提供了程序間傳送資料塊的方法，每個資料塊都可以被認為是有一個型別，接受者接受的資料塊可以有不同的型別；我們可以通過傳送訊息來避免命名管道的同步和阻塞問題；訊息佇列與命名管道一樣，每個資料塊都有一個最大長度的限制；我們可以將每個資料塊當作是一

風河虛擬化元件使用說明（15）—— 為Windows系統安裝VNIC驅動及共享記憶體驅動(on target)

參考Guest Guide文件"Windows VNIC Driver Overview"和"Windows Shared Memory Driver Overview" 首先將Windows 10作為GuestOS啟動：  進入WindShare網站（http://wind

c/c++ linux 程序間通訊系列4，使用共享記憶體

linux 程序間通訊系列4，使用共享記憶體 1，建立共享記憶體，用到的函式shmget, shmat, shmdt 函式名 功能描述 shmget 建立共享記憶體，返回pic key

檔案記憶體對映mmap解決大檔案快速讀寫問題和程序間共享記憶體

mmap函式主要用途有三個： 1、將一個普通檔案對映到記憶體中，通常在需要對檔案進行頻繁讀寫時使用，這樣用記憶體讀寫取代I/O讀寫，以獲得較高的效能； 2、將特殊檔案進行匿名記憶體對映，可以為關聯程序提供共享記憶體空間； 3、為無關聯的程序提供共享記憶體空間，一般也是將一個普通檔案對映到

程序間通訊（三）共享記憶體

概念： 共享記憶體區是最快的IPC形式。⼀一旦這樣的記憶體對映到共享它的程序的地址空間，這些程序間資料傳遞不再 涉及到核心，換句話說是程序不再通過執⾏行進⼊入核心的系統調⽤用來傳遞彼此的資料。 共享記憶體中的函式： shmget函式： 功能：⽤用來建立共享記憶體 原型

Linux：程序間通訊（匿名管道命名管道）（共享記憶體，訊息佇列，訊號量）

目錄 程序間通訊的介紹 管道 匿名管道 原理： 程式碼實現 匿名管道特性 實現管道符 |  命名管道 命名管道特性 程式碼實現 管道讀寫規則 作業系統中ipc的相關命令 共享記憶體（重點） 生命週期： 程式碼實現 程式碼實現獲

作業系統（11）程序--程序通訊：訊號、管道、訊息佇列、共享記憶體

文章目錄 1. 程序通訊相關概念 1. 通訊流程、屬性、鏈路 2. 程序通訊方式：直接通訊、間接通訊 2. 程序通訊的機制 1. 訊號 2. 管道 3. 訊息佇列

ORA-04031 無法分配 12519000 位元組的共享記憶體 large pool ， unknown obje

分享一下我老師大神的人工智慧教程！零基礎，通俗易懂！http://blog.csdn.net/jiangjunshow 也歡迎大家轉載本篇文章。分享知識，造福人民，實現我們中華民族偉大復興！        

C++：通過C++程式碼簡單理解程序間的通訊機制：共享記憶體

下面用共享對映檔案的方式實現程序間通訊，程式碼可以執行。 一、淺理解 每個程序有自己獨立的空間，一個程序無法訪問其他程序的資料。就好像兩個是互不干涉的個體，想讓它們進行通訊（交換資料），就必須有一段它們都可以訪問到的空間，作為中間介質。在計算機中，可以存放資料的地方分為記憶體和硬

共享記憶體及其用mmap實現共享記憶體

一、什麼是共享記憶體 顧名思義，共享記憶體就是允許兩個不相關的程序訪問同一個邏輯記憶體。共享記憶體是在兩個正在執行的程序之間共享和傳遞資料的一種非常有效的方式。不同程序之間共享的記憶體通常安排為同一段實體記憶體。程序可以將同一段共享記憶體連線到它們自己的地址空間中，所有程序都可以訪問共享記

Linux（高階程式設計）8————程序間通訊4（共享記憶體）

共享記憶體是什麼？ 因為程序之間是相互獨立的，他們有各自程序地址空間，那麼他們需要通訊時就要藉助核心來為他們建立橋樑，像之前我們瞭解的管道、訊息佇列就是核心做的工作來為程序間通訊架的橋樑。共享記憶體也是核心為程序間通訊駕的一座橋樑，只不過這座橋樑比其他橋樑更優，共享記憶體是核心為需要通訊