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

看板C_and_CPP (C/C++)作者a5000ml (咖啡裡的海洋藍)時間16年前 (2008/10/22 19:24)推噓5(5推 0噓 0→)

留言5則, 5人參與討論串1/1

※ [本文轉錄自 VideoCard 看板] 作者: a5000ml (咖啡裡的海洋藍) 看板: VideoCard 標題: [分享] CUDA 程式設計(6) -- 記憶體時間: Wed Oct 22 19:17:12 2008 因為學弟們覺得需要練習, 所以從這單元開始加入習題, 並把一些原來在本文中的內容移到習題中 ~ ^_^= 第六章記憶體種類 ◆ 簡介 ◆ 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)