[分享] CUDA 程式設計(6) -- 記憶體

[a5000ml (咖啡裡的海洋藍)]

因為學弟們覺得需要練習, 所以從這單元開始加入習題, 並把一些原來在本文中的內容
移到習題中 ~ ^_^=

第六章 記憶體種類

◆ 簡介 ◆

CUDA 的記憶體種類很繁雜, 依重要順序排列, 可分為下六種

        (1) 暫存器      register
        (2) 全域記憶體  global memory
        (3) 共享記憶體  shared memory
        (5) 常數記憶體  constant memory
        (4) 材質快取    texture cache
        (6) 區域記憶體  local memory

其中最重要的前 4 種, 基本的功能如下

        (1) 暫存器     : 執行緒使用到的一般變數 (快速, 預設, 執行緒內部)
        (2) 共享記憶體 : 同一區塊內的執行緒共用 (快速, 區塊內交換資料用)
        (3) 常數記憶體 : 存放整個程式共用的常數 (快速, 有快取, 全域共用)
        (4) 全域記憶體 : 顯示卡中的 DRAM        (很慢, 無快取, 全域共用)

材質快取是獨立於其它記憶體的存取機制, GPU 使用上要透過 API, 但是它有快取,
即使沒有做什麼複雜的最佳化, 都可以達到不錯的效能, 我們將另闢一個單元討論它.
最後的區域記憶體是被動產生的, 我們無法在 CUDA 的 C 程式語法中主動控制它
(除非嵌入組合語言 ptx), 而且它對效能的影響是負面的, 要避免它的出現.



◆ 延伸標籤 ◆

共享記憶體、全域記憶體、常數記憶體三者具有延伸標籤, 我們可以使用這些標籤直接
宣告變數, 並在 GPU 核心中使用傳統 C/C++ 的方式進行一般操作, 取值或取址.

        (1) 全域記憶體  __device__     在檔案範圍宣告
        (2) 常數記憶體  __constant__   在檔案範圍宣告
        (3) 共享記憶體  __shared__     在函式範圍宣告

在主機的存取上, 用這些標籤宣告的變數不能直接存取, 必需先經過 API 進行位址的
解析, 或直接使用 API 進行存取。(詳見操作特性 7)


範例 6.1: 以下程式碼會將陣列 p[] 週期性地填成 0,1,2,0,1,2,0,1,2,0,1,2,....

        __device__ int a[3]={0,1,2};  //在 file scope 宣告

        __global__ void set_array(int* p, int n){
                int j=blockIdx.x*blockDim.x+threadIdx.x;
                int m=gridDim.x*blockDim.x;
                for(int k=j; k<n; k+=m){
                        p[k]=a[k%3];  //週期性地填入
                }
        }

範例 6.2: 同樣的程式使用常數快取來加速

        __constant__ int a[3]={0,1,2};  //在 file scope 宣告

        __global__ void set_array(int* p, int n){
                int j=blockIdx.x*blockDim.x+threadIdx.x;
                int m=gridDim.x*blockDim.x;
                for(int k=j; k<n; k+=m){
                        p[k]=a[k%3];
                }
        }

在兩個範例中的指位器 p 指向的是全域記憶體, 位在顯示卡的 DRAM 上.



◆ 共享記憶體 ◆

共享記憶體特殊之處在於它的共用範圍是整個區塊, 會產生這種設計的原因在於效能　
和硬體成本上的考量, 其用途很廣

        (1) 執行緒之間快速資料交換. (基本功能)
        (2) 存放會多次使用資料, 當做可控制的快取.
        (3) 存放區塊中的共用變數, 減少暫存器的使用量.
        (4) 做為全域記憶體的合併讀取緩衝 (coalesced read)
        (5) 做為資料順序動態重整單元, 避免區域記憶體的介入.

使用時必需配合 __syncthreads() 對執行緒做同步化, 以確保資料已存妥, 避免後續
其它執行緒因不同步而讀取到尚未寫入的記憶體位址, 造成讀取上的錯誤。


範例 6.3: 一維陣列的區域平均 (單一區塊), 未使用共享記憶體

          #define BLOCK_DIM 10

          __global__ void local_average_1(float *r, float *a){
                int j=threadIdx.x;                //j 為執行緒索引

                if(j==0){
                        r[j]=(2*a[j]+a[j+1])/4;         //左邊界=0
                }
                else if(j==BLOCK_DIM-1){
                        r[j]=(a[j-1]+2*a[j])/4;         //右邊界=0
                }
                else{
                        r[j]=(a[j-1]+2*a[j]+a[j+1])/4;  //輸出加權平均
                }
          }

          local_average_1<<<1,BLOCK_DIM>>> (r, a);


          Note: 它的缺點主要在於要計算一個 r[j], 全域記憶體要讀取 3 次,  
                a[j-1], a[j], a[j+1], 在邊界點上要做條件分支, 且分支裡面
                帶的程式碼延遲很重.


範例 6.4: 一維陣列的區域平均 (單一區塊), 使用共享記憶體

          #define BLOCK_DIM 10

          __global__ void local_average_2(float *r, float *a){
                int j=threadIdx.x;                //j 為執行緒索引

                __shared__ float s[BLOCK_DIM+2];  //宣告共享記憶體

                s[j+1]=a[j];      //多執行緒一起將資料載入共享記憶體
                                  //使用 +1 的偏移, 0 和 BLOCK_DIM+1
                                  //兩點做為陣列邊界

                if(j==0){         //只用一個執行緒設定邊界值
                        s[0]=s[BLOCK_DIM+1]=0;
                }

                __syncthreads();  //同步化, 確保資料己存好

                r[j]=(s[j]+2*s[j+1]+s[j+2])/4;  //輸出加權平均
          }


          local_average_2<<<1,BLOCK_DIM>>> (r, a);


    使用亂數測試的結果為 
    +------+------+------+------+------+------+------+------+------+------+
    | 7    | 8    | 6    | 3    | 3    | 9    | 8    | 5    | 9    | 7    |
    +------+------+------+------+------+------+------+------+------+------+
    | 5.5  | 7.25 | 5.75 | 3.75 | 4.5  | 7.25 | 7.5  | 6.75 | 7.5  | 5.75 |
    +------+------+------+------+------+------+------+------+------+------+

    Note: 其中變數重用度為 3, 亦即要計算 1 個 r[j], 全域記憶體只要讀取 1 次, 
          因為計算很少, 效能主要限制於記憶體, 所以加速差不多是 3x.




◆ 操作特性 ◆


(1) 暫存器是所有記憶體中最快的, 執行緒中大部份的局部變數都預設使用暫存器,
    包括陣列(array)也是, 但有些情況下, 它會被編譯器以較慢的區域記憶體取代,
    這些情況包括在執行緒中「同時佔用」過多變數, 以致於使用的暫存器數目
    超過編譯器的限制 (可使用 --maxrregcount N 選項來限制, 預設 N=128 ),
    或是使用動態變數做為索引存取陣列 (因為此時需要引入陣列的順序結構).

    Note:「同時佔用」是指在某一時間點上, 在後續還會被用到的變數才稱為佔用,
          因為編譯器會自動最佳化, 暫存器的使用是自動調配的, 程式碼中宣告的
          變數個數並不等於使用的暫存器數 (單一變數並非直接對應於單一暫存器),
          可使用 --ptxas-options=-v 選項觀察, 會發現有時候明明宣告很多變數,
          結果暫存器使用到的卻很少, 或是稍微調整一下變數使用的次序就會改變
          暫存器使用數目.

(2) 共享記憶體的應用範圍是區塊, 只有同一區塊裡的執行緒才可以共用它, 使用上要
    注意它的大小限制, 以及存取前後要對執行緒同步化, 避免因資料讀寫的先後順序
    錯誤而導致不可預期的資料錯亂, 其效能僅次於暫存器, 是最佳化的一個重點項目.
    另外, 我們無法在設計時期對它進行初始化, 必需等到核心執行時期才能設定它,
    而且也無法在主機中對它進行存取, 現階段 CUDA 只提供 API 指定其大小.

(3) 區域記憶體是暫存器不夠用的時候, 編譯器自動將資料置換到全域記憶體的產物,
    有點像作業系統的頁置換 (page swap), 它對效能的影響是負面的, 而且非常地
    難以捉摸, 所以在最佳化程式的時候, 時常要用編譯器選項 --ptxas-options=-v
    來追蹤它, 深怕一不小心它就蹦出來, 但有時候為了加大區塊中的執行緒數目
    (blockDim), 必需使用編譯器選項 --maxrregcount N 來限制執行緒最大暫存器
    使用量, 沒有它又無法達成這種限制 (同時佔用的變數就這麼多, 一定要置換出去),
    兩者之間往往必需進行妥協, 或使用共享記憶體進行手動置換.

(4) 全域記憶體除了使用 __device__ 標籤宣告, 另外直接透過 cudaMalloc() 等 API
    直接配置的記憶體也算,「全域」顧名思義就是所有執行單元都可以對它進行操作,
    所以凡是放在顯示卡的 DRAM 中, 能夠被所有執行緒操作 (包括在不同區塊中的
    執行緒) 皆稱為全域記憶體. 它讀取寫入是沒有經過快取的, 跟材質快取在硬體上
    屬於不同的 port, 必需配合「合併讀取」(coalesced read, 也就是半個 warp 的
    執行緒同時讀取記憶體中的連續區塊, 使記憶體控制器做一次性的合併發出) 如此
    才能增進其效能 (約 5~10x), 這部份將在後面單元中詳細介紹.

(5) 常數記憶體雖然和材質快取屬於同一個層次, 但因其大小受到限制, 使快取失誤率
    非常低, 所以在執行時期除了第一次使用需要載入時間外, 之後使用和共享記憶體
    一樣快, 它在核心的存取是唯讀的, 只能在檔案中使用初始值的方法設定, 或是在
    主機中透過 API 進行存取, 使用範圍是全域性的.

(6) 材質快取因為有快取做為緩衝, 所以讀取上不需要做合併, 但也因此比直接存取
    全域記憶體稍微慢一點 (數個到數十個週期) , 但仍比「未合併讀取」全域記憶體
    快上甚多, 所以如果在「合併讀取」很複雜的情況下, 使用材質快取是不錯的選擇.
    材質快取是唯讀的, 在快取的區域性上, 除了傳統微處理機的 1D 快取模式外,
    因繪圖需求的緣故, CUDA 亦提供 2D 和 3D 的材質快取, 使用範圍是全域性的,
    這部份也在後面單元中詳細介紹.

(7) 所有的標籤 __shared__, __device__, __constant__ 宣告的變數所對應的位址
    只有在核心中能直接使用, CUDA 將這些變數稱為 Symbol, 在主機中不能直接以
    C/C++ 原生的方式處理 (進行取值或取址), 必需透過 API

    a. __device__ 必需先透過 cudaGetSymbolAddress() 取得位址, 然後才可以呼叫
       其它的主機 API, 例如 cudaMemcpy() 等或丟給其它 kernel 進行操作, 取得的
       位址和 cudaMalloc() 所配置的位址地位相等, 而且可在主機中進行位址偏移
       (offset, 也就是 ptr+offset 這種動作), 它對映的是顯示卡中實體的位址,
       另外亦可透過和 __constant__ 同樣的方式存取.

    b. __constant__ 只能透過 cudaMemcpyToSymbol() 和 cudaMemcpyFromSymbol()
       進行存取. (note: cudaGetSymbolAddress() 不能用)

    c. __shared__ 主機無法直接存取, 只能設定其大小.


(8) 這些記憶體擁有各自的位址空間, 而且載入/儲存的指令不一樣 (see ptx code),
    所以指位器不能混用 (例如無法使用動態的條件來切換同一個指位器, 讓它可以
    在一些條件下指向共享記憶體, 而在其它條件下指向全域記憶體, 這會造成編譯器
    無法由本文關係解析指位器對應的是何種記憶體, 對後續程式進行編碼).



◆ 特性表 ◆

以下是這些記憶體的特性表, 資料 compute 1.0~1.3 通用

+-----------+--------------+--------+-------------+----------+--------------+
|種類       | 延伸標籤     | 生命期 | 存取範圍    | 主機存取 | 硬體配置     |
+-----------+--------------+--------+-------------+----------+--------------+
|暫存器     | (無)         | 區塊   | 執行緒 R/W  |          | on chip      |
|區域記憶體 | (無)         | 區塊   | 執行緒 R/W  |          | DRAM         |
+-----------+--------------+--------+-------------+----------+--------------+
|共享記憶體 | __shared__   | 區塊   | 區塊 R/W    | size     | on chip      |
+-----------+--------------+--------+-------------+----------+--------------+
|材質快取   | (無)         | 程式   | 全域 R(API) | R/W (API)| DRAM + cache |
|常數記憶體 | __constant__ | 程式   | 全域 R      | R/W (API)| DRAM + cache |
|全域記憶體 | __device__   | 程式   | 全域 R/W    | R/W (API)| DRAM         |
+-----------+--------------+--------+-------------+----------+--------------+

             表 6.1 記憶體種類 (按照存取範圍劃分)



+-----------+------------------+--------------------------------+
|種類       | 存取時間(clocks) | 影響效能因素                   |
+-----------+------------------+--------------------------------+
|暫存器     | immediate        |                                |
|共享記憶體 | 4                | 記憶庫衝突 (16KB/MP, 16 banks) |
+-----------+------------------+--------------------------------+
|常數記憶體 | 4,(失誤)400~600  | 快取失誤 (cache 8KB/MP)        |
|材質快取   | 4,(失誤)400~600  | 快取失誤 (cache 6~8KB/MP)      |
+-----------+------------------+--------------------------------+
|區域記憶體 | 400~600          | 不易控制 (compiler auto)       |
|全域記憶體 | 400~600          | 記憶庫衝突, 未合併讀取         |
+-----------+------------------+--------------------------------+

             表 6.2 記憶體存取速度 (按照存取速度劃分)



+--------------+------------+----------+----------+--------+-----------+
| 標籤         | 大小限制   | 使用範圍 | 變數種類 | 初始值 | 外部變數  |
+--------------+------------+----------+----------+--------+-----------+
| __shared__   | 16KB/block | block    | static   | X      | no extern |
| __device__   |            | file     |          | O      | no extern |
| __constant__ | 64KB       | file     | static   | O      | no extern |
+--------------+------------+----------+----------+--------+-----------+

             表 6.3 標籤特性



◆ 練習題 ◆

練習 6.1: 將範例 6.3 和 6.4 改為多區塊的版本, 測試大陣列並比較效能。

練習 6.2: 試用共享記憶體做矩陣乘法和 transpose，並比較未使用共享記憶體的效能。

練習 6.3: 使用 cudaMemcpyFromSymbol() 和 cudaMemcpyToSymbol() 在執行時期修改
          範例 6.1 和 6.2 的 __device__ 和 __constant__ 陣列的內容, 並使用
          cudaGetSymbolAddress() 和 cudaMemcpy() 做同樣的事情, __constant__
          都能成功嗎？若否, 請使用 cudaGetErrorString() 秀出錯誤訊息。

練習 6.4: 使用編譯器選項 --maxrregcount N 和 --ptxas-options=-v 觀察和調整
          練習 6.2 的暫存器與其它記憶體的使用量, 並測試對程式效能的影響.

練習 6.5: 測試使用動態的條件切換同一個指位器, 讓它可以在某些條件下指向
          共享記憶體, 在其它條件下指向全域記憶體, 例如以下兩段單一區塊
          的程式碼, 執行的結果何者正確？

          -------------------------------------------------------
          函式功能: (使用單一區塊, 部份平方) k=0~99
                r[k] = (k<n) ? a[k] : a[k]*a[k];
          -------------------------------------------------------
          __global__ void part_square_1(int *r, int *a, int n){
                int j=threadIdx.x;      //j 為執行緒索引
                __shared__ int s[100];  //宣告共享記憶體

                s[j]=a[j]*a[j];  //先將每個數的平方存入共享記憶體中
                __syncthreads(); //同步化, 確保資料己存好

                int* p;          //使用共同指位器 (混合指向)
                if(j<n){
                        p=a+j;   //指向全域記憶體
                }
                else{
                        p=s+j;   //指向共享記憶體
                }
                r[j]=*p;         //指位器取值
          }


          -------------------------------------------------------
          __global__ void part_square_2(int *r, int *a, int n){
                int j=threadIdx.x;
                __shared__ int s[100];

                s[j]=a[j]*a[j];
                __syncthreads();

                if(j<n){
                        int* p;  //使用不同指位器
                        p=a+j;   //指向全域記憶體
                        r[j]=*p;
                }
                else{
                        int* p;  //使用不同指位器
                        p=s+j;   //指向共享記憶體
                        r[j]=*p;
                }
          }

          -------------------------------------------------------
          主機呼叫方式

          part_square_1<<<1,100>>>(r,a,10);
          part_square_2<<<1,100>>>(r,a,10);
          -------------------------------------------------------


--
※ 發信站: 批踢踢實業坊(ptt.cc)
◆ From: 114.45.208.109
※ 編輯: a5000ml         來自: 114.45.208.109       (10/22 19:22)
推 dkfum:快M
※ a5000ml:轉錄至看板 C_and_CPP                                    10/22 19:24
推 autoupdate:看不懂
推 f7258:快推!! 以免大家說我看不懂...
※ uf2000uf:轉錄至某隱形看板                                       10/22 20:55
推 ninjaxin:看不懂直接按End XDD
推 b24333666:M起來先慢慢看
推 lavatar:推一下
推 suzukihiro:先推.....晚點來研究...........
推 sdk:還有練習題XD...這..這...還蠻完整的 推~
→ sdk:a大該不會之後還要出ptx的教學吧~~~XD
推 linjrming:感謝 研究中 不過沒有環境可以實作 Orz
推 shadowmere:快推!! 以免大家說我看不懂.
推 CDavid:推～
推 Dissipate:XD
→ a5000ml:嗯~~ 也許下次用程式來解釋會比較清楚
→ a5000ml:s大說的 ptx 可能要很後面, 而且它對應的不是真的指令集
→ a5000ml:比較像 java 的虛擬碼, 用 ptx 寫的效能不見得比較快
推 yayax:推~晚點來看~~
推 VictorTom:有看有推:)
推 wowtiger:ptx暫時沒辦法套進去吧 除非能夠搞到PTX所需的key
推 sdk:樓上所說的key是...???
推 vixen:真想知道僱用原作者寫一隻CUDA程式要開價多少,實在太強了!
推 moonshaped:超清楚 利害阿~~~
推 yhd:謝謝你的教學文件，受益良多
→ phkmalloc:範例二, constant memory不會加速 因為不是same addr