[分享] CUDA 程式設計(14) -- 新年來開始 warp 吧

看板C_and_CPP (C/C++)作者a5000ml (咖啡裡的海洋藍)時間16年前 (2009/01/04 10:42)推噓1(1推 0噓 0→)

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

※ [本文轉錄自 VideoCard 看板] 作者: a5000ml (咖啡裡的海洋藍) 看板: VideoCard 標題: [分享] CUDA 程式設計(14) -- 新年來開始 warp 吧時間: Sun Jan 4 10:39:18 2009 親愛的朋友們, 新年快樂!! 感謝過去一年熱情支持與協助的朋友們, 提供國網光碟 DVD 壓縮、FTP、BT 下載, 以及許許多多的資訊, 並且校正我在硬體認識上的謬誤, 使小弟獲益良多. 在這新的一年, 我們邁入 CUDA 進階討論, 其它基本 CUDA 程設, 會再慢慢補足, 比較急的話不妨先看 CUDA ToolKit 的 programming guide, 也可到 google 上搜尋「CUDA 教學」, 有很多很讚的部落格, 他們整理得比我好, 凌亂的生活態度, 想不到竟反應在我寫的文章上, 真是傷腦筋...XD 希望在這新的一年, 能擴大對 GPGPU 這方面計算感興趣的朋友陣容, 能空出時間在 ATI GPGPU 嘗試 programming 與 Apple OpenCL, 不論是進行研究 or 切磋, 期待中文化文件的發展, 以及在 PTT 上成立正式 GPGPU 版, 讓大家能有更好的交流空間. Yours Truly, 咖啡裡的海洋藍 2009 ◆ 何謂 warps ============================================================================= 現階段 warp 指 32 個執行緒, 在硬體上為實際的 SIMT 群組, SM (串流多處理器, Stream Multiprocessor) 的 8 個執行單元, 輪流執行這 32 個執行緒, 達成在 4 個 clocks 之下完成整個 warp 操作 (一般的運算指令, 如果是 global memory 的存取, 當然會需要更多的時間, 但 warp 會在 issue 完指令後, 把工作交給記憶體管線處理, 然後 free to work on 其它的 warp). ◆ Grid/Block/Warp/Thread 架構 ============================================================================= 我們提過 __syncthreads() 對效能很負面, 因為一旦使用這個指令, multi-threads 就無法隱藏 memory access latency, 而這卻是 SIMT 比起傳統設計最大的優勢之一, 解決的辦法之一就是讓 SM 可以同時執行數個 blocks, 使得部份 block 裡的 threads hang 住時, SM 仍可執行其它的 block 來隱藏 latency. 一般人通常認為一個 SM 只能執行一個 block, 這是錯誤的觀念, 事實上一個 SM 具有同時執行數個 blocks 的能力, 但設定上也要做更多調整, 畢竟天下沒有白吃的午餐, 為了簡化 software 設計起見, 初學者並不建議在這方面涉入太深, 但進階最佳化時, 這卻是基本的(programming guide 都有寫), 不過仍屬於控制 blocks 階段. 然而, 絕少文獻提及, SIMT 的群組模型必需要視為 grid/block/warp/thread 架構, 大部份都只提 grid/block/thread (包括 programming guide), 甚至 CUDA 實作時, warps 也被排除在正規軟體架構之外, 只是視為是 hardware 的一部份, 理解上也被當作是 optional 的. (早期的 CUDA 版本中, warp 並不以明顯的方式呈現, 直到 1.2 版後, 才開始有 warp vote 等指令出現) ◆ 對 warp 進行精確控制的好處 ============================================================================= 可曾想過我們可以不用 __syncthreads(), 就可對 shared memory 做存取嗎? Well... 這的確是可以做到的, 因為 warp 是 hardware locking 的同步執行群組, 本身就 synchronized 了. 也就是透過對 warps 的精確控制, 我們可以在標準的 grid/block/thread 架構下, 再細切一層 warp, 在這層中不需要 __syncthreads(). 例如在 SDK 的 scan 這支程式中, 實做 tree reduction 時必需 __syncthreads(), 但我們可以在 warp 層先做不用 __syncthreads() 的 reduction, 而且因為 warp size 是固定, 使用 32 個 threads 來 tree reduce 64 筆 data, 可以直接寫入 6 個指令, 然後再把 global 的那套搬到 block 層再搞一次, 就可以有效的去除許多 __syncthreads(), 而且各 warps 也可以非同步執行. 另一個好處是, 我們在 global memory access 之後, 到 __syncthreads() 之間, 利用 warp 的這種特性可插入許多指令, 更能有效的緩衝 latency hidden 的壓力, 使得 SIMT 的效果呈現得更好. 再來就是, 我們以前時常用到的老招數, 應用在某些地方只用到少量 threads 時, 例如使用單一 thread 載入數個邊界點, 可以分散在數個 warp 中, 避免只使用一個 thread 造成的 I/O serial 現象, 有效的運用 SM 中的多個記憶體單元. 還有很多應用不勝枚舉, 小弟也還在探索中... ◆ warp 的執行緒定址 ============================================================================= 如果指定 blockDim = 512, 則它包含了 16 個 warps, 執行緒定址如下 warp 0: tid=0~31 warp 1: tid=32~63 warp 2: tid=64~95 ... warp 15: tid=480~511 當指定的 blockDim 不是 32 的整數倍時, 最後的 warp 不完整, 硬體上仍佔一個 warp, 只是多餘的單元被 disable 掉而己, 例如 blockDim = 100 包含了 4 個 warps warp 0: tid=0~31 warp 1: tid=32~63 warp 2: tid=64~95 warp 3: tid=96~100 (incomplete, 101~127 disabled) 為了方便操作, 我通常會把 blockDim 設成 2D/3D 的型式, CUDA 的 lowest dimention 是由 x 算起, 所以 threadIdx.(y,z) 會形成 warp 的 2D index, 例如 blockDim=Dim3(32, 5, 1) -- 2D 的方式, 配置 5 個 warp 的陣列 warp 0: threadIdx.y=0 warp 1: threadIdx.y=1 ... warp 5: threadIdx.y=5 blockDim=Dim3(32, 3, 4) -- 3D 的方式, 配置 3x4 的 warp 陣列 warp (0,0): threadIdx.y=0 threadIdx.z=0 warp (0,1): threadIdx.y=0 threadIdx.z=1 warp (0,2): threadIdx.y=0 threadIdx.z=2 warp (0,3): threadIdx.y=0 threadIdx.z=3 ... warp (2,0): threadIdx.y=2 threadIdx.z=0 warp (2,1): threadIdx.y=2 threadIdx.z=1 warp (2,2): threadIdx.y=2 threadIdx.z=2 warp (2,3): threadIdx.y=2 threadIdx.z=3 ◆ 硬體執行配置 1 (compute 規格) ============================================================================= 一個 SM 可同時執行數個 blocks, 其限制在於 register 和 shared memory 的大小, 以及一些規格上的限制 (see compute 版本的規格), 例如在 1.0 版中規定 max # of active blocks per SM = 8 max # of active warps per SM = 24 max # of active threads per SM = 768 這裡的 active 指的是 SM 當下正在執行它, 也就是已經 load 到 SM 中執行了, 其中第 3 條 rule 比較沒有義意, 因為由第 2 條已決定 (24*32=768), 且當存在 incomplete warps 時, 通常是 meet 不到的, 但在純 software 的觀點上 (不理會 warp 時), 有時候會覺得它反而比較容易理解. Anyway, 這些 rules 是用來避免 GPU 配置過多執行緒給 SM 的硬性限制, 否則當 shared memory = 0 時, 所有的 blocks 全都發給同一個 SM 執行, 那還得了. ps. 在支援 compute 1.2 的硬體上, 對這些規格有些異動 max # of active blocks per SM = 8 max # of active warps per SM = 32 max # of active threads per SM = 1024 ◆ 硬體執行配置 2 (shared memory 的影響) ============================================================================= 再來就是 shared memory 限制, 它是實作在每個 SM 上的, 現階段總量只有 16KB, 未使用或 8 blocks 的使用量不滿 16KB 時, SM 就會按照 default 規格來發配 warps, 否則 GPU 會切割出最大的 block 數來配置, 例如 shared = 3KB 時 16KB / 3KB = 5 所以 1 SM 會配置 5 blocks, 這些 blocks 雖然在同一個 SM 中執行, 但為了維持 Grid/Block 架構的彈性, 它們仍是無法再透過 shared memory 彼此溝通的. 當 shared > 8KB (超過 SM 總量的一半) 時, 一個 SM 就只能執行一個 block. 另外通常系統會使用到一些 shared memory, 所以要用 --ptxasoptions=-v 選項來看實際的 shared memory 使用量. ◆ 硬體執行配置 3 (register 的影響) ============================================================================= 在 registers 的用量方面, 影響的不只 compile 過不過, 也影響 SM 對 block 的配置, 考慮如下問題: G200 有 64KB register space, 若 1 個 block 使用 registers 大於 SM 總量之半 (例如 50KB), 假設不使用任何的 shared memory, 且每個 block 使用的 threads 和 warps 很少 (使它 bypass 前面的限制, 例如 blockDim = 20), 難道 GPU 要配給 SM 8 blocks 嗎? 這樣一來, 只有其中一個 block 的 register 能夠 in core, 其它的 7 blocks 呢? 所以在硬體設計上, 讓 register space 對 SM 分配的 block 進行限制是必要的. ◆ 如何 fulfill max # of active threads per SM ============================================================================= 也許你會說： God damn... @#$^%, 這麼多限制條件, 有什麼辦法讓 SM 能開啟全部 threads? 管它是 1.0 版的 768 還是 1.2 版的 1024 threads. Well...這是可行的: (1) 首先在 block 配置時必需要是 complete warp (blockDim = 32 的整數倍), 不然 1024 個 threads 一定有些會被 disabled. (2) 再來要理解 block 數不必完全填滿, 不然在 shared memory 控制上壓力會很大, 例如在 1.2 版中, 寧願配 2 個 blockDim=512, 也不願 8 個 blockDim=128, 起碼前者可用的 shared 有到 8 KB, 後者只有 2KB, 當然這要看應用而定. (3) 限制 register 和 shared memory 的使用量. (4) 在 launch 時使用 device query 的 API 來決定 blockDim 大小, 讓程式對後續版本具有相容性. (optional) 例如向量加法就很容易這樣做, 因為它使用的 register 很少且 shared memory = 0 compute 1.0/1.1 版: blockDim=384 -> 2 blocks, full 768 threads compute 1.2/1.3 版: blockDim=512 -> 2 blocks, full 1024 threads 也可以配置 blockDim=256 (768,1024 的最大公因數, 8 warps) compute 1.0/1.1 版: blockDim=256 -> 3 blocks, full 768 threads compute 1.2/1.3 版: blockDim=256 -> 4 blocks, full 1024 threads note: 配置多點的 block 在 SM 上較可隱藏 __syncthreads() 的負面效果, 所以要對 shared memory 使用量進行妥協. -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 114.45.219.173 -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 114.45.219.173