2016年3月19日 星期六

閱讀筆記:GPU的指令運算細節

最近看到了一篇 Nervana 發表的文章
他們宣稱可以讓 CUDA 矩陣乘法運算速度在 Maxwell 上達到理論上限的 90% 以上
裡面除了用到 register blocking(從別的 paper 借來用的技巧)
也講解了很多一般不太可能知道的 NVIDIA GPU 運作細節
讓人不禁懷疑作者是不是從 NVIDIA 的相關部門拿到資料
知道這些細節之後,就可能產生出高效率的矩陣乘法 CUDA kernel

Latency

在 pipeline 的系統中一定有 latency
就是一個指令下去之後有多久才會真正反應在 register 裡面
例如說 Maxwell GPU 裡面實行基本浮點數或是整數運算之後
要過 6 個 cycles 才會被寫到 register 裡面(我記得從 Tesla 就是這個數字了)
而如果是 predicator 的話比較久,要 13 cycles
另外有一些延遲是非固定值
例如 special register 像是 threadIdx, blockIdx 等等要 20 個 cycles 左右
記憶體的話甚至更慢 shared 要 30+,global 則要 100+ 個 cycles

此外,有時候一個指令要過一陣子才會從 register 讀值
例如說 load/store 進入 pipeline 之後
延遲 2(shared)/4(global) 個 cycle 才會從 register 拿要存取的記憶體位置
大約 20 cycles 才會拿完(因為他離 register 比較遠?)
這段時間內都不能動到這些 register
要不然會算出錯誤的記憶體位置

ILP/Superscalar

因為 pipeline latency 所以不可能每個 cycle 都有指令可以下去
例如說
c=a+b
d=e+c

如果第一個指令之後下一個 cycle 馬上下第二個指令
那麼就會讀到錯誤的 c,因為 c 還沒被寫進去(GPU 要 6 cycles)

CPU 上有兩種解法
一個是在不影響正確性的情形下
插入不相關的指令
c=a+b
p=q+r
s=t+u
...
d=e+c

或者找不到不相關指令的時候
乾脆不做事 (no oepration, NOP)
c=a+b
NOP*5
d=e+c

另外相反地,如果有多個硬體
也有可能在 1 cycle 內同時跑多個指令
這個就是一般說的 ILP (instruction level parallelism)
c=a+b, p=q+r
...
d=e+c

一般可以由編譯器來決定這些東西
不過有很多情況是不是編譯時期決定的
這時候硬體也必須有一定的處理方式
CPU 常見的作法有 scoreboard 跟 Tomasulo (register renaming)
現在的 CPU 上面應該都是後者的變形
而 GPU 應該偏向前者

另外,在 GPU 上面因為同時執行了一定數量的 thread/warp
所以也可以透過切換 warp 的方式來隱藏 latency
c=a+b (warp 1)
c=a+b (warp 2)
c=a+b (warp 3)
...
d=e+c (warp 1)

當然,前面可以使用的 ILP 技巧也可以使用
而且除了同一個 warp 可以並行之外
GPU 多了讓不同 warp 在同一個時間執行的可能性
c=a+b (warp 1), c=a+b (warp 2)
c=a+b (warp 1), p=q+r (warp 1)

Maxwell 上的 Scheduling

為了省電的考量(應該吧)
NVIDIA Kelper 之後的架構在執行檔中插入了額外的資訊
可以用 cuobjdump -sass 看出來
                                          /* 0x003fb400e3a007e6 */
MOV R1, c[0x0][0x20];                     /* 0x4c98078000870001 */
S2R R0, SR_TID.X;                         /* 0xf0c8000002170000 */
SET.GE.AND P0, PT, R0, c[0x0][0x140], PT; /* 0x4b6d038005070007 */

圖中紅色的部份就是額外資訊
額外的資訊大小是在 Kelper 中是程式的 1/7,Maxwell 是 1/3
不過一般來說多出來的資訊不可能看出來是什麼
但是 Nervanas 的人就是能知道他的意義
而且文章也用這個資訊推論出一些 NVIDIA GPU 的架構

Maxwell 中額外的資訊每 21 bits 分配給一個指令
其中分成兩大部份

前 4 bits 表示 register reuse
應該跟 ISCA 或是 MICRO 2013 前後的 register file cache, operand cache
這兩篇論文有很大的關聯(忘記詳細是哪些會議了)
每個 operand slot 都有兩個 cache entry
而這些 bit 就是提示說「應該把這個 register 放進去 cache」
這樣作可以節省從 register file 抓資料的電力

後面的 17 bits 表示這個指令 schedule 相關的資訊
從高位到低位又分成五項

  1. wait barrier (6)
  2. set read barrier (3)
  3. set write barrier (3)
  4. yield (1)
  5. stall count (4)

Stall Count

stall count 就是這個指令要等幾個 clock 才能放行下一個指令
像是這種常見的 dependency 又沒插入額外指令的話就是
c=a+b (stall = 6)
d=e+c

0 也是可以允許的值,這種情形就是前述的 ILP
在 white paper 是叫做 dual issue
這時 schduler 可以同使把這個指令跟下面 schedule 下去
不過文章中說似乎只允許不同硬體單元的 dual issue
例如 memory 跟 ALU 同時 issue
c=a+b (stall = 0)
p=q+r

另外某些指令像是 branch, syncthread, return
其 stall count 必須大於 5,要不然後面的指令會偷跑

yield

yield 的話,如果有寫過 python 應該有看過這個東西
程式執行到那邊的時候會把執行權轉交出去
CUDA 裡面就是把這個 warp 的執行權轉給下一個 warp
例如說這個案例
c=a+b (stall = 1)
p=q+r
因為都沒有遇到 stall 超過 1 的情形
所以 scheduler 似乎會優先跑同一個 warp
c=a+b (warp 1)
p=q+r (warp 1)
c=a+b (warp 2)
p=q+r (warp 2)

有 yield 的情況下 warp 跑完後執行權會轉交給下一個 warp
作者說加上這個以後程式會執行得比較平均
c=a+b (warp 1, yield)
c=a+b (warp 2, yield)
p=q+r (warp 1)
p=q+r (warp 2)

作者說有一個奇怪的點是
如果 stall 超過 12 的指令不加上 yield 會出錯

Barrier

Maxwell 的架構中有 6 個 barrier
六個中每個 barrier 都可以設成 read barrier 或是 write barrier 其一
是為了對付不定長度的 latency 使用的
例如說 memory 或是 special register (blockIdx, threadIdx...)
真正把資料寫入 register 的時間不一樣
有時甚至超過 4 bits stall count 可以表示的數字

write dependency 是最好理解的,用來避免 RAW hazard
就是不希望資料還沒寫進 register 就被讀取了就可以 set write dependency
例如說有一個指令是 "load a byte 0x0000 to register 15"
我可以 "set barrier 3 as write barrier"
之後用到 register 15 的第一個指令就必須 "wait barrier 3"
這樣這個指令就會等 barrier 3 變成 resolved 才會被 scheduler 丟下去執行

read dependency 的需則是因為某些指令不會把 register 複製起來計算
(主要是 memory operation)
memory operation 大約要花 20 cycles 才能把需要的位置算出來
如果有人這時候寫入記憶體的話就會算錯位置
read barrier 大致上跟 write barrier 一樣,只是 stall 的時間不同
write barrier 是 resolved 之前不能讀取 register
read barrier 是 resolved 之前不能寫入 register
也就是說他是拿來免除 WAR hazard 用的

至於 WAW 因為 GPU 不會偷把後面的指令拿來執行 (out of order issuing, OOO)
所以不需要處理

這樣就可以知道 wait barrier 為什麼是 6 bits
而 read/write 是 3 bits

wait barrier 就是指定的要等哪些 barrier
一個指令可以等前面的多的指令的完成,所以是 bit flag
總共要 6 bits

而一個指令可以對 read write 個別設定一個 barrier
0-5 就是要設定哪個
magic number 7 是不設 barrier
所以是 3 bits

作者說有一個要注意的點是
如果 set barrier 要在緊接的下一個指令使用的話
要等一個 cycle,也就是 stall count = 2
因為 barrier 必須要 1 cycle 之後才會生效

雜七雜八

文章中還有一些不知道歸類成什麼的細節參數
既然文章中有了就順便列出來

第一,register file 有分成數個 bank
Kelper 跟 Maxwell 都是四個
如果一個指令同時存取多個相同 bank 的話
會導致 throughput 下降
看 operands 有幾次衝突,就會下降成 1/2 甚至 1/3 (ex: fused mult-add)

第二,這個作者也找出了 memory request queue 大約的長度
shared memory 大約是 ~50/SM
global memory request 的話 ~8/warp
texture memory request 的話更少,似乎是 1/warp

小結

看完這篇文章之後長了不少知識
不過還是無法想像作者怎麼推敲出這些資訊的

沒有留言:

張貼留言