GPU指令集技術(shù)分析
本文將兩篇文章整理了一下。
參考文章鏈接如下：
https://zhuanlan.zhihu.com/p/391238629
https://zhuanlan.zhihu.com/p/166180054
一．GPGPU- 指令執(zhí)行設(shè)計
本節(jié)主要內(nèi)容：
? GPGPU指令執(zhí)行簡介
o GPGPU指令執(zhí)行流水線
o GPGPU指令執(zhí)行吞吐的影響因素
o GPGPU指令執(zhí)行的特點
? 指令設(shè)計中的一些原則與思路
o 指令長度的問題
o 指令集設(shè)計與ILP的一些相關(guān)性
o 復(fù)合操作與附加操作
o 立即數(shù)操作數(shù)和Constant Memory操作數(shù)
o 格式的通用性與信息具體化
? 簡單聊聊一些具體的指令
o FMA, MUL, ADD 系列
o IMAD, LEA, IADD3
o LOP3
o MUFU
o FSETP
o LDG/STG, LDS/STS, LDL/STL
o BAR
? 結(jié)語
二．CUDA微架構(gòu)與指令集-指令發(fā)射與warp調(diào)度
CUDA指令的發(fā)射和warp調(diào)度問題。
? 指令發(fā)射的基本邏輯，主要是指令發(fā)射需要滿足的條件，幾代架構(gòu)發(fā)射指令的一些簡單描述等等。
? control codes，主要是它與指令發(fā)射和warp調(diào)度的關(guān)系。
? warp Scheduler的功能。
? 峰值算力的計算方法和達成條件。
現(xiàn)在分開來敘述
一．GPGPU- 指令執(zhí)行設(shè)計
GPGPU指令吞吐和指令集設(shè)計的一些問題。NV GPU的機器碼指令集叫SASS。
指令集是微架構(gòu)與用戶對接的途徑。指令集相當于硬件提供給軟件的API（或者也可以認為指令集是輸入前端，微架構(gòu)是執(zhí)行后端）。按照
GPGPU指令執(zhí)行簡介
GPGPU指令執(zhí)行流水線
首先，先簡單介紹一下通用處理器的指令執(zhí)行邏輯。對RISC V一種簡單的五級流水線實現(xiàn)的描述：

1. Instruction fetch cycle (IF)：主要是獲得Program Counter對應(yīng)的指令內(nèi)容。
2. Instruction decode/register fetch cycle (ID)：解碼指令，同時讀取輸入寄存器的內(nèi)容。由于RISC類指令的GPR編碼位置相對固定，所以可以直接在解碼時去讀取GPR的值。
3. Execution/effective address cycle (EX)：這是指令進入執(zhí)行單元執(zhí)行的過程。比如計算memory地址的具體位置（把base和offset加起來），執(zhí)行輸入輸出都是GPR或輸入含立即數(shù)的ALU指令，或者是確認條件跳轉(zhuǎn)指令的條件是否為真等。
4. Memory access (MEM)：對于load指令，讀取相應(yīng)的內(nèi)存內(nèi)容。對于store指令，將相應(yīng)GPR的值寫入到內(nèi)存地址處。RISC類指令集通常都是load-store machine，ALU指令不能直接用內(nèi)存地址做操作數(shù)（只能用GPR或立即數(shù)），因而通常ALU指令沒有memory access。
5. Write-back cycle (WB)：將相應(yīng)的ALU計算結(jié)果或memory load結(jié)果寫入到相應(yīng)的GPR中。
   與復(fù)雜的CISC指令集相比，多數(shù)GPGPU的指令集還是比較接近load-store machine，總體來說與RISC更相似一些。GPGPU典型微架構(gòu)可以簡單表示為下圖：
   

GPGPU指令執(zhí)行流程示意圖
這個圖是微架構(gòu)的模塊示意圖，而非流水線示意圖。實際執(zhí)行流程中Fetch、Decode、Execution這三步是必須的，而Memory access顯然只針對memory指令，write-back則只針對需要寫回的指令（比如memory load，帶GPR輸出的指令等）。
注：流水線的配置與ALU的latency有很大的關(guān)系。比如Volta前FFMA的延遲是6cycle，Volta及之后FFMA的延遲是4cycle，這絕對與流水線的改進有關(guān)。不過，這里的Latency并不是所有流水線的級數(shù)。因為Latency在程序中的表現(xiàn)形式是：一個指令發(fā)射后，其結(jié)果需要多少周期才能就緒（也就是能被其他指令使用）。兩個back-to-back dependent的ALU指令（比如FFMA R0, R1, R2, R0; FFMA R0, R3, R4, R0;），前一個FFMA只要在第二個FFMA讀取操作數(shù)之前把結(jié)果寫回GRF，那后一個FFMA就可以得到正確值。對應(yīng)到上面的5級流水線形式，就是前一個指令的WB要在后一個指令的ID前執(zhí)行完就行（相當于4cycle延遲），最開始的IF那一級是不影響的。CPU對于這種形式的依賴還有更激進的旁路邏輯（forwarding），可以直接在前一個ALU的EX后把結(jié)果直接送給后一個ALU的EX當輸入，從而減少流水線的bubble，提高性能。NV的GPU應(yīng)該是沒有這么緊湊的forwarding，但是NV的operand collector可以作為一個公共的操作數(shù)中轉(zhuǎn)站，理論上前一個ALU的結(jié)果寫回到operand collector就可以被下一個ALU看到了，不一定要回到GRF。
由于GPU運行模型的復(fù)雜性，在Decode后Execution前，還有大量其它的步驟：比如scoreboard的判斷（主要用來保證每個指令的執(zhí)行次序，防hazard），warp Scheduler的仲裁（多個eligible warp中選一個），dispatch unit和dispatch port的仲裁（發(fā)射或執(zhí)行資源沖突時需要等待），還可能有讀取作為立即數(shù)操作數(shù)的constant memory，讀取predicate的值生成執(zhí)行mask，等等。在執(zhí)行過程中，也有很多中間步驟，比如輸入GPR操作數(shù)的bank仲裁，跳轉(zhuǎn)指令要根據(jù)跳轉(zhuǎn)目標自動判斷divergence并生成對應(yīng)的mask，訪存指令要根據(jù)地址分布做相應(yīng)的請求廣播、合并等等。在寫回這一級的時候，由于一些指令的完成是異步的（比如一些內(nèi)存指令），所以也可能需要GPR端口的仲裁，等等。
當然，步驟雖然多而瑣碎，但未必都會新增單獨的一級流水。GPU應(yīng)該是為了簡化設(shè)計和節(jié)約功耗，不愿意把流水線拉得太細長，因而很多操作都是塞在同一級流水線里，各種組合邏輯非常復(fù)雜。這樣也導(dǎo)致它的主頻往往就不能太高。比如最近幾代NV GPU旗艦和次旗艦的主頻：

可以看到主頻基本都在1~2 GHz之間，次旗艦的頻率往往比旗艦要稍高一些（這里選的公版頻率，但非公也大致是這個趨勢），有些低端芯片頻率可能還會更高一點。而如今（2021年）常見的桌面端x86 CPU，基準頻率3~4 GHz，最大睿頻4~5 GHz是很尋常的事。很多Arm CPU的大核，主頻也能接近甚至超過3 GHz。當然，這么比也許不是特別公平。因為多數(shù)獨立GPU的功耗很大，會極大的限制頻率提升。一些眾核CPU的頻率也會比少核版的降一些，不過差別不會太大（類似上面GPU的旗艦與次旗艦的關(guān)系）。但即使算上這些，GPU的主頻比常見CPU的主頻還是顯著低一些（實際上帶核芯顯卡的CPU里，GPU頻率往往也是顯著小于CPU頻率）。這里面具體的因果關(guān)系我也不是特別明白，感覺肯定還是有些故事的~
GPGPU指令執(zhí)行吞吐的影響因素
指令執(zhí)行吞吐一般指的是每個時鐘周期內(nèi)可以執(zhí)行的指令數(shù)目，不同指令的吞吐會有所不同。通常GPU的指令吞吐用每個SM每周期可以執(zhí)行多少指令來計量。對于多數(shù)算術(shù)邏輯指令而言，指令執(zhí)行吞吐只與SM內(nèi)的單元有關(guān)，整個GPU的吞吐就是每個SM的吞吐乘以SM的數(shù)目。而GPU的FMA指令（通常以F32計算）往往具有最高的指令吞吐，其他指令吞吐可能與FMA吞吐一樣，或是只有一半、四分之一等等。所以很多英文文檔會說FMA這種是full throughput，一半吞吐的是half rate，四分之一的是quarter rate等。當然，有些微架構(gòu)下也會有1/3、1/6之類非2的冪次的比率。NV GPU近幾代微架構(gòu)的常見指令吞吐如下：

CUDA算術(shù)邏輯指令吞吐表
從圖中可以發(fā)現(xiàn)，指令吞吐不僅與指令類型有關(guān)，還與微架構(gòu)具體設(shè)計實現(xiàn)有關(guān)。主要會受到以下一些因素的影響：

1. 功能單元的數(shù)目。絕大多數(shù)指令的功能都需要專用或共享的硬件資源去實現(xiàn)，設(shè)計上配置的功能單元多，指令執(zhí)行的吞吐才可能大。顯然，只有最常用的那些指令，才能得到最充分的硬件資源。而為了節(jié)約面積，很多指令的功能單元會相互共享，所以他們的吞吐往往也會趨于一致。比如浮點的FFMA、FMUL都要用到一個至少24bit的整數(shù)乘法器（32bit浮點數(shù)有23bit尾數(shù)，小數(shù)點前還有1bit）。以前一些處理器有24bit的整數(shù)乘法指令，兩者乘法器就可以共用，從而具有相同的吞吐（不過NV最近幾代好像都沒有這個指令，ptx以及內(nèi)置函數(shù)的24bit乘法應(yīng)該是多個指令模擬的）。而FADD雖然用不上那個乘法器，但可以與FFMA共用那個很寬的加法器，以及一些通用的浮點操作（特殊數(shù)的處理，subnormal flush之類）。32bit的整數(shù)乘法因為需要更寬的乘法器，有的就不會做成full throughput，甚至可能被拆分成多個指令（比如Maxwell和Pascal用三個16bit乘法指令XMAD完成一次32bit整數(shù)乘法）。Turing的IMAD應(yīng)該是有意識的加寬了，所以32bit的IMAD與FFMA吞吐一樣，但印象中帶64bit加數(shù)的IMAD應(yīng)該還是一半。再比如一些超越函數(shù)指令（MUFU類，比如rcp，rsq，sin，exp之類），由于實際使用量相對不會太頻繁，多數(shù)是1/4的throughput。
2. 指令Dispatch Port和Dispatch Unit的吞吐。這個在之前的專欄文章也詳細講過。一個warp的指令要發(fā)射，首先要eligible，也就是不要因為各種原因stall，比如指令cache miss，constant immediate的miss，scoreboard未就位，主動設(shè)置了stall count等等。其次要被warp scheduler選中，由Dispatch Unit發(fā)送到相應(yīng)的Dispatch Port上去。Kepler、Maxwell和Pascal是一個Warp Scheduler有兩個Dispatch Unit，所以每cycle最多可以發(fā)射兩個指令，也就是雙發(fā)射。而Turing、Ampere每個Warp Scheduler只有一個Dispatch Unit，沒有雙發(fā)射，那每個周期就最多只能發(fā)一個指令。但是Kepler、Maxwell和Pascal都是一個Scheduler帶32個單元（這里指full-throughput的單元），每周期都可以發(fā)新的warp。而Turing、Ampere是一個Scheduler帶16個單元，每個指令要發(fā)兩cycle，從而空出另一個cycle給別的指令用。最后要求Dispatch Port或其他資源不被占用，port被占的原因可能是前一個指令的執(zhí)行吞吐小于發(fā)射吞吐，導(dǎo)致要Dispatch多次，比如Turing的兩個FFMA至少要stall 2cycle，LDG之類的指令至少是4cycle。更詳細的介紹大家可以參考之前的專欄文章。
3. GPR讀寫吞吐。絕大部分的指令都要涉及GPR的讀寫，由于Register File每個bank每個cycle的吞吐是有限的（一般是32bit），如果一個指令讀取的GPR過多或是GPR之間有bank conflict，都會導(dǎo)致指令吞吐受影響。GPR的吞吐設(shè)計是影響指令發(fā)射的重要原因之一，有的時候甚至占主導(dǎo)地位，功能單元的數(shù)目配置會根據(jù)它和指令集功能的設(shè)計來定。比如NV常用的配置是4個Bank，每個bank每個周期可以輸出一個32bit的GPR。這樣FFMA這種指令就是3輸入1輸出，在沒有bank conflict的時候可以一個cycle讀完。其他如DFMA、HFMA2指令也會根據(jù)實際的輸入輸出需求，進行功能單元的配置。
4. 很多指令有replay的邏輯，這就意味著有的指令一次發(fā)射可能不夠。這并不是之前提過的由于功能單元少，而連續(xù)占用多輪dispath port，而是指令處理的邏輯上有需要分批或是多次處理的部分。比如constant memory做立即數(shù)時的cache miss，memory load時的地址分散，shared memory的bank conflict，atomic的地址conflict，甚至是普通的cache miss或是TLB的miss之類。根據(jù)上面Greg的介紹，Maxwell之前，這些replay都是在warp scheduler里做的，maxwell開始將它們下放到了各級功能單元，從而節(jié)約最上層的發(fā)射吞吐。不過，只要有replay，相應(yīng)dispath port的占用應(yīng)該是必然的，這樣同類指令的總發(fā)射和執(zhí)行吞吐自然也就會受影響。
   幾個需要注意的點：
5. 指令發(fā)射吞吐和執(zhí)行吞吐有時會有所區(qū)別。有些指令有專門的Queue做相應(yīng)的緩存，這樣指令發(fā)射的吞吐會大于執(zhí)行的吞吐。這類指令通常需要訪問競爭性資源，比較典型的是各種訪存指令。但也有一些ALU指令，比如我們之前提過的Turing的I2F只有1/4的吞吐，但是可以每cycle連發(fā)（也就是只stall 1cycle）。不過多數(shù)ALU指令的發(fā)射吞吐和執(zhí)行吞吐是匹配的。
6. 要注意區(qū)分指令吞吐與常說的FLOPS或是IOPS的區(qū)別。通常的FLOPS和IOPS是按乘法和加法操作次數(shù)計算，這樣FMUL、FADD是一個FLOP，FFMA是兩個FLOP。這也是通常計算峰值FLOPS時乘2的由來。但是有些指令，可以計算更多FLOP。比如HFMA2 R0, R1, R2, R3;可以同時算兩組F16的FMA，相當于每個GPR上下兩個16bit分開獨立計算（類似于CPU的SIMD指令），所以SM86以前的架構(gòu)HFMA2的指令吞吐與FFMA是一樣的，只是每條指令算4個F16的FLOP，而FFMA是2個F32的FLOP。這也就是TensorCore出現(xiàn)前F16的峰值通常是F32兩倍的原因。DFMA由于輸入寬度比FFMA再翻倍，所以功能單元做成一半就能把GPR吞吐用滿（這里說的是滿配Tesla卡，消費卡F64常有縮減）。因此，在TensorCore出現(xiàn)以前，通常的Tesla卡HFMA2、FFMA的指令吞吐一樣，DFMA吞吐是一半，而看峰值FLOP就是H:F:D=4:2:1的關(guān)系。TensorCore出現(xiàn)后，指令（比如HMMA）本身的吞吐和指令入口的GPR輸入量沒有變化，但由于同一個warp的指令可以相互共享操作數(shù)數(shù)據(jù)，一個指令能算的FLOP更多了，因而峰值又提高了。當然，這里說的是一般情況，實際上根據(jù)產(chǎn)品市場定位的不同，有些功能可能會有所調(diào)整。
7. SM86（Ampere的GTX 30系列）的F32比較另類。Turing把普通ALU和FFMA（包括FFMA、FMUL、FADD、IMAD等）的PIPE分開，從而一般ALU指令可以與FFMA從不同的Dispatch Port發(fā)射，客觀上是增加了指令并行度。NVIDIA對CUDA Core的定義是F32核心的個數(shù)，所以Turing的一個SM是64個Core。Ampere則把一般ALU PIPE中再加了一組F32單元，相當于一個SM有了128個F32單元（CUDA Core），但是只有64個INT32單元。也就是說SM86的F32類指令的吞吐是128/SM/cycle，但其中有一半要與INT32的64/SM/cycle共享。或者說，Turing的F32和INT32可以同時達到峰值（包括A100），而SM86的INT32和F32不能同時達到峰值。
   GPGPU指令執(zhí)行的特點
   與傳統(tǒng)的x86 CPU相比，GPGPU在指令執(zhí)行的邏輯上有很多獨特的地方。
   靜態(tài)資源分配：GPU有一個很重要的設(shè)計邏輯是盡量減少硬件需要動態(tài)判斷的部分。GPU的每個線程和block運行所需的資源盡量在編譯期就確定好，在每個block運行開始前就分配完成（Block是GPU進行運行資源分配的單元，也是計算Occupancy的基礎(chǔ)）。典型的運行資源有GPR和shared memory。GPU程序運行過程中，一般也不會申請和釋放內(nèi)存（當然，現(xiàn)在有device runtime可以在kernel內(nèi)malloc和free，供Dynamic Parallelism用，但這個不影響當前kernel能用的資源）。CPU在運行過程中有很多所需的資源是動態(tài)調(diào)度的。比如，x86由于繼承了祖上編碼的限制，ISA的GPR數(shù)目往往比物理GPR少，導(dǎo)致常常出現(xiàn)資源沖突造成假依賴。實際運行過程中，通常會有register renaming將這些ISA GPR映射到不同的物理GPR，從而減少依賴（有興趣的同學可以研究下tomasulo算法）。GPU沒有這種動態(tài)映射邏輯，每個線程的GPR將一一映射到物理GPR。由于每個線程能用的GPR通常較多，加上編譯器的指令調(diào)度優(yōu)化，這種假依賴對性能的影響通常可以降到很低的程度。
   每個block在運行前還會分配相應(yīng)的shared memory，這也是靜態(tài)的。這里需要明確的是，每個block的shared memory包括兩部分，寫kernel時固定長度的靜態(tài)shared memory，以及啟動kernel時才指定大小的動態(tài)shared memory。雖然這里也分動靜態(tài)，但指的是編譯期是否確定大小，在運行時總大小在kernel啟動時已經(jīng)確定了，kernel運行過程中是不能改變的。
   其實block還有一些靜態(tài)資源，比如用來做block同步的barrier，每個block最多可以有16個。我暫時沒測試到barrier的數(shù)目對Occupancy的影響，也許每個block都可以用16個。另一種是Turing后才出現(xiàn)的warp內(nèi)的標量寄存器Uniform Register，每個warp 63個+恒零的URZ。因為每個warp都可以分配到足額，應(yīng)該對Occupancy也沒有影響。另外每個線程有7個predicate，每個warp有7個Uniform predicate，這些也是足額，也不影響Occupancy。
   GPU里還有一種半靜態(tài)的stack資源，通常也可以認為是thread private memory或者叫l(wèi)ocal memory。多數(shù)情況下每個線程會用多少local memory也是確定的。不過，如果出現(xiàn)一些把local memory當stack使用的復(fù)雜遞歸操作，可能造成local memory的大小在編譯期未知。這種情況編譯器會報warning，但是也能運行。不過local memory有最大尺寸限制，當前是每個線程最多512KB。
   順序執(zhí)行：亂序執(zhí)行是CPU提高CPI的一個重要途徑，但亂序執(zhí)行無論是設(shè)計復(fù)雜度還是運行控制的開銷都很大。CPU的亂序執(zhí)行可以把一些不相關(guān)的任務(wù)提前（相關(guān)的也可以亂序，但要求順序提交），從而提高指令并行度，降低延遲。而GPU主要通過Warp切換的邏輯保持功能單元的吞吐處于高效利用狀態(tài)，這樣總體性能對單個warp內(nèi)是否stall就不太敏感。
   雖然GPU一般是順序執(zhí)行，但指令之間不相互依賴的時候，可以連續(xù)發(fā)射而不用等待前一條指令完成。在理想的情況下，一個warp就可以把指令吞吐用滿。當然，實際程序還是會不可避免出現(xiàn)stall（比如branch），這時就需要靠TLP來隱藏這部分延遲。
   顯式解決依賴：既然是順序執(zhí)行，但同時又可以連續(xù)發(fā)射，那怎么保證不出現(xiàn)數(shù)據(jù)冒險呢？NV GPU現(xiàn)在主要有兩類方式：第一種是固定latency的指令，通過調(diào)節(jié)control codes中的stall count，或者插入其他無關(guān)指令，保證下一條相關(guān)指令發(fā)射前其輸入已經(jīng)就位；第二種是不固定latency的指令，就需要通過顯式的設(shè)置和等待scoreboard來保證結(jié)果已經(jīng)可用。在x86的CPU中，memory結(jié)果的可見性是通過緩存的一致性來控制的，這樣read-after-write之類的組合可以通過cache的可見性來保證，但多線程的情況也需要通過coherence和memory consistency model來保證。GPU本身運行就是多線程的，同一個warp內(nèi)也是通過scoreboard來保證次序。但多個warp之間，GPU也需要維護相應(yīng)的coherence和memory consistency model，具體大家可以參考PTX文檔: Memory Consistency Model。
   當然這個邏輯雖然是這么設(shè)計的，估計偶爾也會有出bug的時候。Maxwell和之前的架構(gòu)偶爾能看見編譯器往程序內(nèi)插一些NOP。大概就是硬件上有問題，靠編譯器來強行修補。Turing上似乎已經(jīng)比較少見了。
   指令設(shè)計中的一些原則與思路
   指令長度的問題
   至少從Kepler開始，SASS就是定長的了。Kepler每條指令64bit，每8條指令含一條control code，之后的Maxwell、Pascal每條指令還是64bit，但是control code變成了每4條指令一條。Volta、Turing、Ampere都是每條指令128bit，每條都自帶control code。
   定長和變長有什么差別呢？一般講體系結(jié)構(gòu)的書上會用RISC與CISC來做對比，因為一般CISC指令集是變長的，比如x86。而RISC則通常是定長的。定長的好處之一是解碼器可以提前解碼，且一般解碼開銷小。因為首先指令等長，每個指令的范圍是確定的，指令定界不依賴于前一個指令的解碼。其次RISC在編碼的時候一般各個域的位置和長度比較整齊，解碼相對說來自然也更簡單一些。而變長則只能順序解碼，不得到前一條指令長度，后一條指令就不知道從哪里開始，當然就無法解碼。同時由于長度是變化的，每個操作數(shù)的類型和內(nèi)容都可能會變化，解碼也就更繁瑣一些。不過變長有一個好處就是可以壓縮一些常用指令的長度，從而減少程序大小。而定長就沒有這種方式。
   不過，現(xiàn)在很多架構(gòu)和指令集設(shè)計都漸漸趨同。比如x86雖然變長（從一個Byte到十幾個Byte），但它解碼前會先定界，然后可以經(jīng)過預(yù)解碼變成一系列的Micro Operation，這樣真正執(zhí)行的時候也類似RISC。而RISC也不都是定長的，比如ARM的Thumb模式可以混用16bit和32bit的指令，RISC-V也可以加長指令實現(xiàn)特定的功能擴展。
   對于GPU來講，我覺得長指令還是有很大的好處的。我的理由如下：
8. 指令Cache是GPU中命中率最高的Cache之一（constant cache也許是真的No.1）。一個warp 32線程已經(jīng)讓指令解碼和調(diào)度之類的開銷大大均攤了。而同時大部分代碼都會被成千上萬個warp運行，這個開銷還可以被平攤得更小。當然，每個指令變長會加重指令Cache的容量和吞吐負擔，但它的格式也可以做得更整齊更有規(guī)律，從而大大簡化解碼過程。像x86這種編碼復(fù)雜的指令集，連解碼都會添加相應(yīng)的cache（Micro-Op Cache）。而只要指令夠長夠整齊，解碼就可以做得夠簡單。這高ICache命中率的GPU來說，總體來講還是賺的。或者你可以認為，x86的Micro-op cache主要利好循環(huán)這種高重用代碼，而GPU每條代碼都與循環(huán)類似，那不如直接就用解碼后的指令做輸入。
9. 前面提過指令執(zhí)行的顆粒度。由于GPU流水線和發(fā)射邏輯的限制，每條指令都有基礎(chǔ)開銷（至少占用一次Dispatch Unit和相應(yīng)Port），那在同一條指令中塞進更多信息，就成為減少指令基礎(chǔ)開銷的重要手段。指令越長，能編碼的信息就越多，比如指令內(nèi)嵌的立即數(shù)就可以支持更長。典型的如ALU的立即數(shù)操作數(shù)，load/store的offset，branch的跳轉(zhuǎn)目標等。同時，多數(shù)指令也可以支持更復(fù)雜的modifier，從而表述能力和可編程性都更強。同時，更整齊規(guī)范的格式也更有利于編譯器做優(yōu)化，性能的可及性更好。
10. 指令夠長，就可以加入更多的調(diào)度和控制信息，從而簡化硬件自主判斷，拓寬指令觀察視野。Control code就是一個典型的例子。大部分的Control code其實是編譯期信息，沒有control code，其中很多依賴的判斷就需要專門的硬件邏輯去檢測和處理。而受限于硬件實現(xiàn)的復(fù)雜度，其視野肯定不如編譯時的全局視野寬。因此，通過這些控制信息，可以更好的利用編譯器在編譯時獲得的先驗知識，減少硬件自主分析和處理，從而簡化硬件設(shè)計，提高面積利用率的同時也簡化設(shè)計和驗證過程。
    Volta開始的128bit指令，我覺得還是一個很有意義的嘗試。也許有人覺得這個也許可以做點內(nèi)存壓縮，不過我感覺意義不是太大，壓縮和解壓其實也算是編碼解碼過程，尋址還會增加負擔，未必便宜。
    指令集設(shè)計與ILP的一些相關(guān)性
    GPU兩個最重要的并行邏輯，ILP（Instruction Level Parallelism）和TLP（Thread Level Parallelism，TLP有時在CPU語境下也代指Task Level Parallelism，兩者還是有所區(qū)別），兩者在隱藏延遲中都有重要作用。ILP的邏輯主要是靠前一條指令不需要執(zhí)行完成就能發(fā)射下一條無關(guān)指令，而TLP則是通過warp之間切換來隱藏延遲。從另一個角度講，ILP和TLP都可以增加可發(fā)射指令的數(shù)目，盡量減少功能單元的閑置，從而提高硬件利用效率。
    ILP是線程內(nèi)（更準確的說是Warp內(nèi)）的并行邏輯，影響ILP的主要因素有兩種，一是指令之間的依賴性，二是指令的資源競爭或沖突。依賴分顯式和隱式。顯式依賴主要是數(shù)據(jù)的相關(guān)性，隱式依賴則與資源競爭很相似，主要是兩個指令都要使用某個特定含義的公共資源。典型的顯式依賴如：
    FFMA R3, R1, R2, R0; // sm_75: stall 4 cycles
    FFMA R6, R4, R5, R3;
    而隱式的依賴比如這種：
    // sm_61
    1: IADD RZ.CC, R0, R1 ; // set condition code as carry
    2: IADD.X R2, RZ, R2 ; // use condition code as carry
    3: IADD RZ.CC, R3, R4 ;
    4: IADD.X R5, RZ, R5 ;
    IADD可以把進位存到專門的CC寄存器（類似x86的carry flag），然后IADD.X可以把這個CC寄存器當成carry讀進來再做加和。由于CC寄存器只有一個，雖然2和3兩條指令沒有數(shù)據(jù)依賴，仍然不能把2和3互換以隱藏1-2和3-4之間的延遲。而在Turing里這種指令已經(jīng)可以顯式的用Predicate來存儲carry，如：
    // sm_75
    1: IADD3 R4, P0, R0, R2, RZ ; // set P0 to carry out
    2: IADD3 R10, P1, R6, R8, RZ ; // set P1 to carry out
    3: IADD3.X R5, R1, R3, RZ, P0, !PT ; // use P0 as carry in
    4: IADD3.X R11, R7, R9, RZ, P1, !PT ; // use P1 as carry in
    這樣1、3和2、4兩組指令就可以interleave，正確性互不影響，從而相互隱藏延遲，提高ILP。
    這種相對更獨立的指令集設(shè)計其實有點類似函數(shù)式編程：操作專門carry寄存器可以看做是stateful的操作，改成可編程的Predicate后就成為只與輸入輸出有關(guān)的stateless操作，不改變機器狀態(tài)。也可以認為是所有需要用carry寄存器做輸入輸出的指令，都需要被序列化。通過與當前機器狀態(tài)解耦，獲得更大的指令調(diào)度自由度，編譯器的后端優(yōu)化也會更加方便。
    NV GPU在很早就開始有意識的淘汰這種含隱式輸入輸出的指令。比如早期使用隱式棧的call、ret、break等等（當然其實這和ILP關(guān)系已經(jīng)不太大了）。例如SM61中要return時需要先顯式的用PRET設(shè)置某個隱式的調(diào)用棧，然后直接用RET返回。顯然在RET時這個棧必須是對應(yīng)當初PRET設(shè)置的值（中間能不能再進出棧沒仔細研究），否則就會出錯。而Turing直接使用帶GPR地址的指令進行操作，就消除了這種隱式棧的操作過程，減少了指令之間復(fù)雜依賴對編譯器的干擾。
    // RET in sm_61
    PRET 0x258 ;
    …
    RET;

// RET.REL in sm_75
MOV R20, 32@lo((Z7argtestPiS_S + .L_9@srel)) ; // relocation with addend
MOV R21, 32@hi((Z7argtestPiS_S + .L_9@srel)) ;
…
RET.REL.NODEC R20 `(Z7argtestPiS_S);

// RET.ABS in sm_75
RET.ABS R32 `(Z7argtestPiS_S);
再稍微擴展一點。x86中有control register來控制如何做浮點數(shù)的rounding，是否做subnormal的flush（x87 FPU control register控制普通的FPU指令，MXCSR控制SSE指令）。但在SASS中，FFMA、FMUL、DFMA等浮點運算指令，每個指令都可以自主控制是否打開flush（使用FTZ modifier），如何做rounding（RM, RP, RZ, RN）。這就意味著每個指令可以自主決定當前指令的運行方式，而不用改變機器狀態(tài)。不同模式混用時，就不需要保存和恢復(fù)control register了。
那再再擴展一點，x86其實也有控制FP exception的寄存器，NV GPU里是怎么操作的呢？我好像沒看見，感覺是被去掉了。這還是一個挺值得思考的問題~
當然，也不是說所有的隱式都應(yīng)該變成顯式。比如每個Warp都有一個隱式的active mask，用來標記當前warp中divergence的情況。active mask與指令predicate的“AND”會共同決定當前指令是否起作用。那把mask交由指令顯式操作有意義嗎？我覺得沒有，因為這沒有太多額外的可編程價值。首先divergence造成的mask變化只能被分支或分支同步指令修改，其他指令需要控制效果直接用predicate就可以了，沒必要操作mask。其次，這個操作本身是非常固化的，增加相關(guān)操作指令并不會帶來新功能，反而會增加指令負擔（相當于每次可能有divergence的分支時，都要顯式的保存和設(shè)置mask，遍歷不同divergence的分支時就更麻煩了）。因此，在有predicate的情況下，active mask還是做成隱式的比較合理。當然，volta后的Independent Thread Scheduling也要自主控制和依賴內(nèi)部的mask狀態(tài)，就更沒法做成顯式的了。不過，雖然mask是隱式的，但divergence后重新converge一般是顯式的（通過BSSY和BSYNC指令），否則程序就不知道應(yīng)該在哪個點join了。
復(fù)合操作與附加操作
前面也提到了每個指令都有基礎(chǔ)開銷，那盡量在一個指令里塞進更多事情以分攤基礎(chǔ)開銷，就成為指令改進的一個重要方向。畢竟每個指令都要占用發(fā)射機會，但提高頻率很困難，也就是發(fā)射的總指令數(shù)有限，那就只好讓每個指令多做事。
比如一般的ALU指令可能有1~3個輸入操作數(shù)，為了盡量利用GPR的讀吞吐和增加每個指令的操作能力，SASS里其實有單指令多操作的情況。最常見的就是FMA（包括FFMA、DFMA、HFMA2和IMAD等），它可以一條指令同時計算乘和加兩個操作：d=a*b+c。Tensor Core的MMA類的指令也通常是3個輸入操作數(shù)。這里以Turing為例，常見的多輸入或多操作指令有：
FFMA R5, R4, R5, R11 ;
DFMA R4, R2, R4, R2 ;
HFMA2 R3, R13, R9, R3 ;
IMAD.WIDE R10, R9, R4, R10 ;
IADD3 R36, -R11, R22, -R34 ;
LOP3.LUT R9, R15, R9, R16, 0xfe, !PT ;

LEA.HI.X R7, R3, R51, R2, 0x3, P0 ; // Load Effective Address: d = (a<<b) +c
SHF.R.U64 R3, R11, R0, R12 ; // Funnel shift
PRMT R61, R60, R58, R61 ; // Byte permute

I2IP.S8.S32.SAT R0, R1, R0, R2 ; // Integer To Integer Conversion and Packing
IDP.4A.S8.S8 R9, R20, R25, R9 ; // Integer Dot Product and Accumulate
IMMA.8816.S8.S8 R36, R50.ROW, R74.COL, R36 ; // Integer Matrix Multiply and Accumulate
HMMA.1688.F32 R0, R184, R200, R0 ; // Half Matrix Multiply and Accumulate
這個列表并不完整，比如ISCADD雖然在Turing指令集里，但是編譯器似乎更傾向于用LEA和IMAD，所以好像很少見到了。還有warp shuffle指令SHFL，也可以接受4個GPR做操作數(shù)，不過這不是通常意義上的ALU指令，其實更接近memory指令一些，所以也沒列出來。然后是大部分非tensor類的整數(shù)和位操作指令都有Uniform datapath的對應(yīng)版，這里也沒列。
這些多操作指令往往都需要更多的操作數(shù)。上一期我們講吞吐的時候也提到了，GPR的吞吐用不滿也就浪費了，所以一般ALU最多是3個輸入GPR。其實每cycle最多能讀4個，不過多少還是要留點余量給其他異步指令（如內(nèi)存讀寫指令），否則搶占GPR的端口和搶占發(fā)射機會是一個效果。立即數(shù)和Predicate因為不占用GPR的吞吐，所以還可以額外加。比如LOP3.LUT R9, R15, R9, R16, 0xfe, !PT ;就是3個GPR+1立即數(shù)+1Predicate的輸入。于是，NV在設(shè)計指令集的時候又經(jīng)常會在一些ALU操作后加一個免費的Predicate操作。比如SETP類的指令本來就是通過比較生成一個predicate，但它也可以順手把生成的predicate與另一個predicate做與、或、異或等，這樣就可以把一些鏈式的bool操作（比如 a < M && b < N && c < K）附帶在比較中。包括前面說的LOP3，以及PLOP3等，都具有這種額外的predicate輸入。
FSETP.NEU.AND P0, PT, RZ, c[0x0][0x16c], PT ;
ISETP.LT.OR P0, PT, R9, R8, P0 ;
ISETP.EQ.XOR P4, PT, R7, R1, P4 ;
還有另外一類額外的附加操作，比如float intrinsic里有一個函數(shù)叫__saturatef(float x)，它會把輸入clamp到[0,1]這個區(qū)間里。但是它并不占用一個指令，因為float指令自帶一個modifier叫.SAT，如FFMA.SAT R6, R2, R7, R0 ; 會自動對結(jié)果R6做saturate，并沒有額外開銷。類似的這種飽和操作在一些整數(shù)操作中也會出現(xiàn)（比如I2I.S16.S32.SAT R11, R11 ;表示把S32的輸入saturate到S16的范圍）。
復(fù)合操作和附加操作可以大大增強指令的表述能力，用更少的指令做更多的事情。但也要看到，它增加了指令的復(fù)雜度，對編譯器后端優(yōu)化也提出了更多的要求。同時，在硬件設(shè)計上它也會有一些額外的開銷。如果這些操作不是特別常用，對指令運行開銷反而是負擔，所以多數(shù)還是需要一些功耗控制措施。比如IADD3其實多數(shù)時候都有一個是RZ，LOP3也通常只有兩個有效輸入。至于SETP后帶Predicate的情況，其實占比也不高。所以這些在具體設(shè)計和實施上，應(yīng)該還是需要做相應(yīng)的優(yōu)化的。
立即數(shù)操作數(shù)和Constant Memory操作數(shù)
由于Volta前的SASS指令只有64bit，一般ALU的立即數(shù)操作數(shù)中只有19bit的編碼。這對于一些特定的指令是不太夠的，所以早期也有特定的32bit立即數(shù)的指令，如FADD32I，FMUL32I，MOV32I等等。不過這些指令最多只有2個輸入操作數(shù)，FFMA這種就沒法弄了。而計算很多函數(shù)值（比如sin）一般都是從高次多項式近似開始，需要計算大量給定系數(shù)的FMA，這些要用立即數(shù)就只能用mov先把32bit立即數(shù)寫入GPR，然后再FFMA。Volta后每個ALU包括FFMA都能用滿32bit立即數(shù)，就沒有這個問題了。
當然，實際Volta前的多項式逼近不是用立即數(shù)實現(xiàn)的，而是用Constant Memory。因為Constant Memory需要的編碼更少（幾bit的bank編碼+16bit的地址編碼），而且Constant memory還可以用在64bit的DFMA上，通用性更強。這也是constant memory對指令表達能力的一個重要貢獻。
Constant memory與立即數(shù)在做操作數(shù)時還有不少區(qū)別： 1. Constant memory運行有overhead，啟動kernel前需要初始化。而立即數(shù)是hardcode在編碼里，沒有額外overhead。 2. 立即數(shù)是編譯期常數(shù)，編譯時就必須知道值。Constant memory可以做運行期常數(shù)，也就是啟動kernel前才需要得到具體值，在運行前是可調(diào)的。 3. Constant memory有容量限制，當前應(yīng)該是每個bank 64KB。立即數(shù)的總?cè)萘縿t是程序能寫多長就可以多大。 4. Constant memory運行期開銷會大一些，因為它畢竟是內(nèi)存，需要相應(yīng)的cache和load單元支持。 5. 立即數(shù)的編碼能力比較有限，當前128bit的指令也就32bit立即數(shù)。而Constant memory只要在指令格式中約定好，32bit和64bit都可以，將來擴展到128bit甚至更多也不是不行。
當前在CUDA程序里，編譯期常數(shù)是放在Constant memory還是做立即數(shù)是編譯器決定的。一般來講，32bit能精確表示的是立即數(shù)（比如float的所有數(shù)，double的1.0，2.5之類），需要64bit才能精確表示的會用constant memory。
注：傳運行期常數(shù)，直接用kernel的參數(shù)不也可以嗎？有什么區(qū)別呢？
其實很相似，但確實也有區(qū)別。因為kernel并沒有什么來自host的調(diào)用棧，kernel的參數(shù)其實也是用constant memory存儲的，在kernel啟動前會由驅(qū)動自動初始化。所以兩者從性質(zhì)上很相似，但又有點差別。kernel參數(shù)的scope是當前kernel，每次啟動kernel理論上都要重新初始化參數(shù)區(qū)的constant memory，這其實也是kernel啟動的overhead之一（不太確定cuda Graph會不會做優(yōu)化）。而用戶自己設(shè)置和copy的constant memory，scope至少是module（可以簡單認為是一個cu文件，或者更具體一點是cubin文件）。這就意味著同一個module的kernel都可見。所以如果同一個module內(nèi)的kernel共用大量參數(shù)，且中間不會更改，那就只需要初始化一次（重復(fù)啟動同一個kernel也是同理）。
這甚至可以是一個極致優(yōu)化的例子。把所有kernel的輸入?yún)?shù)都放在某個struct里，然后把struct復(fù)制到constant memory，就可以不用參數(shù)啟動這所有的kernel了。這應(yīng)該會節(jié)省一點overhead（如果cuda graph不能做優(yōu)化）。
PS：其實這個操作的意義主要在于接口的可擴展性。需要增加新參數(shù)的時候，只需要在struct里新加一個成員變量就行，不用去改函數(shù)參數(shù)和調(diào)用kernel的地方了。相信我，這個有時候真是很方便！你當然也可以在kernel參數(shù)里傳一個struct指針，這樣代碼量也相當。但是你需要自己把struct內(nèi)容memcpy到device上，讀參數(shù)時還會增加內(nèi)存負擔。因為constant在SM內(nèi)有專門的cache，load時又有broadcast機制，hit時幾乎沒有額外開銷。
格式的通用性與信息具體化
很多指令可以支持大量的modifier，把一些輸入信息更加具體化。這里我們舉一個典型的Turing IMAD的例子：
IMAD R4, R5, R0, -R4 ;
IMAD.HI.U32 R49, R5, c[0x0][0x1d8], RZ ;
IMAD.WIDE.U32 R20, R17, c[0x0][0x168], R20 ;

IMAD.MOV.U32 R31, RZ, RZ, c[0x0][0xc] ;
IMAD.SHL.U32 R0, R2, 0x8, RZ ;
IMAD.IADD R2, R7, 0x1, R11 ;
前3個其實都沒啥特別。IMAD本身是算整數(shù)的d=a*b+c，類似浮點的FMA（前面也提到了，實際上Turing的IMAD和FFMA共用了功能單元）。單獨的IMAD就是三個32bit的計算，保留低32bit。IMAD.HI就是保留高32bit，IMAD.WIDE表示加數(shù)c是64bit（a、b還是32bit），輸出d也是64bit。這都算是常規(guī)操作。
后三個就有點意思了，IMAD.MOV其實就相當于MOV，因為它的輸入a、b肯定都是RZ（SASS中的恒零寄存器），所以輸出d肯定等于c。用IMAD做MOV主要是MOV指令和普通INT32是一個Dispatch Port，而IMAD是FMA的port，兩者錯開有利于提高ILP。但我直接寫IMAD.U32 R0, RZ, RZ, R1; 與 IMAD.MOV.U32 R0, RZ, RZ, R1; 有什么區(qū)別呢？功能上兩者應(yīng)該是一樣的，但功耗上可能不太一樣。因為IMAD會用到一個很長的乘法器，功耗會很大。但如果事先知道這個乘法器不用了，那硬件上就可以繞過或是用一些處理不觸發(fā)它，從而節(jié)省功耗。后面兩個指令也是類似，IMAD.SHL相當于b一定是2的冪次，IMAD.IADD相當于b一定是1，其實也是不需要完整的乘法器功能。不寫SHL或IADD不影響結(jié)果，但顯式的告訴功能單元，可以讓功能單元得到更明確具體的信息，從而進行一些優(yōu)化。
那能不能讓硬件自己檢測呢？當然可以，但是未必劃算。首先，如果大多數(shù)IMAD指令都是不帶特殊性的操作數(shù)，那為了檢測這種case的功耗優(yōu)化方案就會給硬件造成額外負擔。就像cache一樣，hit的時候自然是好，但加了cache往往也會劣化miss時的開銷。如果miss比率太高的話，加cache就是個負優(yōu)化。同樣，在這里如果大部分情況執(zhí)行路徑可以被優(yōu)化，那自然可以省功耗。但如果多數(shù)情況其實優(yōu)化不了，那就白白損失了優(yōu)化檢測的這部分硬件開銷，變成了負優(yōu)化。當然，增加這個modifier本身也是要看使用頻率的，如果極少用到，那其實也是降低了面積效率。這都需要對使用場景和規(guī)劃做準確的判斷。
這個邏輯其實也和control code有些相通之處，盡量把編譯期的信息融入到程序中，讓程序盡量明確運行邏輯。這其實也許體現(xiàn)了NV硬件設(shè)計中一個更上層的指導(dǎo)思想：軟件盡量具體地告訴硬件運行方式，而硬件則盡量減少自身的判斷和復(fù)雜監(jiān)測邏輯，無腦的運行軟件的指示即可。這樣硬件設(shè)計盡量簡化，不但能節(jié)約面積開銷，測試、驗證復(fù)雜度都能降低。當然，這不是一件簡單的事情，想有效的安排和傳遞各種編譯信息，肯定需要軟硬件協(xié)同設(shè)計。這對軟硬件的架構(gòu)規(guī)劃和協(xié)作，包括對應(yīng)用、編譯器、指令集、硬件架構(gòu)等統(tǒng)籌規(guī)劃，還是提出了很高的要求，很體現(xiàn)功力。
這里IMAD做這些優(yōu)化應(yīng)該是和編譯器有確定配合的，編譯器在合適的時候會優(yōu)先選擇這些模式。其實這些功能也可以有其他一些實現(xiàn)方式，但是多則惑，也沒有必要。開銷大致等價的實現(xiàn)中選一個就可以了。
一些具體的指令
指令集內(nèi)容實在太多了，有很多東西我沒有靠譜的輸入，瞎猜也意思不大。我就簡單列一些SASS中我覺得還比較有意思的點，供大家參考。這里基本都是Turing的指令。
FMA,MUL, ADD 系列
FFMA是衡量GPU算力的標桿之一。從指令吞吐角度講，FFMA一般都是最高的那一組。FFMA功能上相當于FMUL+FADD（FFMA精度更好，因為只做了一次舍入），那FMUL與FADD與FFMA共用單元就很合理。不過另一個有意思的地方是，FADD的reuse cache用的slot是1、3，是不是暗示它其實是FFMA的套殼，只是把FFMA第二個操作數(shù)固定成1？當然，更省的辦法應(yīng)該是有直接的短路邏輯。
DFMA, DMUL, DADD系列具有不定長的延遲，需要靠scoreboard去控制依賴。我覺得一是因為FP64相關(guān)單元配置變化很大，有的還是多個SMSP之間share，可能存在競爭，這樣延遲不可控。二是它的吞吐可能很小（有的卡是1/16），指令延遲可能會超過stall count能有效表示的范圍，所以還是用scoreboard比較靠譜。
這些都支持32bit立即數(shù)做操作數(shù)。D系列雖然每個操作數(shù)是64bit，但如果后32bit都是0，那也可以放進32bit立即數(shù)里，用的時候做padding就可以了（指數(shù)位還是按FP64來，只是截了尾數(shù)后32bit）。HFMA2、HMUL2、HADD2則相當于把它當成兩個16bit的立即數(shù)。
IMAD, LEA, IADD3
IMAD在前面已經(jīng)介紹過一些了。這里可以稍微再補充一些相關(guān)的。
IMAD.MOV.U32 R31, RZ, RZ, c[0x0][0xc] ;
IMAD.SHL.U32 R0, R2, 0x8, RZ ;
IMAD.IADD R2, R7, 0x1, R11 ;

IMAD R5, R5, c[0x0][0xc], RZ ;
IMAD.X R21, RZ, RZ, R7, P0 ;
前三個我們已經(jīng)介紹過，為了控制乘法器的開銷，有相應(yīng)的modifier去明確輸入。但是可以發(fā)現(xiàn)，加數(shù)為RZ時沒看到專門的modifier。IADD3其實多數(shù)情況下也只有兩個操作數(shù)，第三個是RZ的概率很大。也沒有看到有modifier去優(yōu)化這個，也許單純檢測RZ開銷不大，當然也可能是加法器的開銷沒那么大，不是很敏感。
另外就是IMAD在接受carry輸入時，有的前兩個輸入也都是RZ，但是這次就沒見到相應(yīng)的modifier。不知道是沒設(shè)置還是沒有，可能是加了carry這個邏輯不適用，也可能是編譯器沒考慮？也不是很清楚。亦或是這種概率很小，可以忽略？
與IMAD的一些功能有共通之處的還有LEA指令。這是計算地址時常用的指令。IMAD是d=ab+c，LEA則特殊一點是d=(a<<b)+c。這樣算X[i]這種地址時，相當于isizeof(type) + *X。如果size是2的冪次，那就可以用LEA。LEA與INT32是一組，IMAD是FMA那組，兩個也可以配合使用。
LOP3
LOP3還是一個挺有意思的指令，它可以完成三個32bit數(shù)的任意按位邏輯運算。其實邏輯也很簡單，把這個映射關(guān)系看成一個函數(shù)d=F(a, b, c)，abcd都是bool值。輸入狀態(tài)有2^3=8種可能，那每個F函數(shù)都可以用這8個輸入時的輸出來完全表示。用8bit的編碼（立即數(shù)查找表，immeLUT）就可以指定這樣一個邏輯函數(shù)。更具體的用法和介紹大家可以參考PTX文檔。

LOP3 ImmeLUT
粗看起來好像挺厲害的，其實多數(shù)時候也沒有什么大用。我曾經(jīng)搜過官方庫里的所有LOP3指令，看到的情況是：絕大部分情況下第三個操作數(shù)是RZ，也就是說3輸入的按位邏輯函數(shù)其實還是挺少見的。如果真的要找場景，可能密碼學或者挖礦之類會有些不錯的應(yīng)用。
不過細品一下，其實還是有一些妙用的。比如我有兩個32bit數(shù)a，b，我需要把a的某些bit和b的bit拼在一起（對位替換），怎么弄呢？這其實就是一個簡單的按位邏輯函數(shù)c[i] ? a[i] : b[i]。這樣lop3就能用上。那什么時候會用到這種功能呢？其實還是有一些。比如copysign，需要把a的符號位替換掉b的符號位。再比如浮點數(shù)里有些指數(shù)、尾數(shù)之類的操作，也是有一些用處的。
更具體一點的例子，比如y= (x>=0) ? 1 : -1，這相當于就是把x的符號位復(fù)制給整數(shù)1。不過…… 這個一條LOP3指令做不到，因為這里c我們已經(jīng)用了立即數(shù)（0x70000000，也就是選中第一個符號位bit），a或b就不能再用立即數(shù)了，等于說這里需要先把1或0x70000000移到GPR里，才能用LOP3來操作。
MUFU
MUFU是SASS中計算各種超越函數(shù)的指令。數(shù)學上，超越函數(shù)是相對代數(shù)函數(shù)（有限次加、減、乘、除、開方等組合）而言。但硬件上不能用多項式表述的好像都歸在超越函數(shù)里了。多項式求值一般就是乘和加，不需要其他的指令。而超越函數(shù)，就需要通過一些其他指令或是軟件逼近來實現(xiàn)。有的地方也叫特殊函數(shù)，我感覺是不太合理，就這些初等函數(shù)怎么也談不上很特殊吧……
常見的MUFU類的指令有：
MUFU.RSQ R5, R10 ;
MUFU.RCP R3, R14 ;
MUFU.EX2 R9, R8 ;
MUFU.LG2 R10, R9 ;
MUFU.COS R9, R19 ;
MUFU.SIN R10, R19 ;

MUFU.RCP64H R3, R7 ;
MUFU.RSQ64H R11, R29 ;
其實數(shù)學函數(shù)求值這個事情還是有很多可聊的點。
首先，這些指令都是給出近似值。對精度要求高的話還是需要調(diào)用相應(yīng)的數(shù)學函數(shù)庫再做軟件實現(xiàn)。一般說來，有的數(shù)學函數(shù)可以迭代更新（比如求倒數(shù)reciprocal，RCP），那就可以從近似值開始用不動點迭代得到更精確的值。另一種是可以做argument的range reduction，把全范圍的參數(shù)縮放到能比較精確計算的范圍內(nèi)。比如對于x具有二進制形式x=a2^e, log2(x) = log2(a2^e)=e + log2(a)。那只需要精確計算log2(a)（其中a在1~2之間）就可以得到精度很高的解，而MUFU.LG2在某些特定范圍內(nèi)是可以滿足精度要求的。還有一種就是sin和cos這種，它沒有迭代形式，也沒有特別好的手段做規(guī)約，那就先規(guī)約到一個不太大的范圍（比如正負pi/4），然后用一個高次多項式去近似（不是泰勒展開，大家有興趣可以去研究一下Remez的minimax方法）。所以MUFU.SIN只是在近似計算中出現(xiàn)，精確計算中可能用不著它。
當然，軟件實現(xiàn)上還是有很多自由度，這里只是給出了一種選擇方式。NV的數(shù)學函數(shù)庫不直接開源，但是有相應(yīng)的LLVM的bitcode（NVVM目錄下的libdevice.*.bc），可讀性還比較強，可以參考。
從這些實現(xiàn)可以看到，大量輔助計算仍然是用加、乘這種初級運算完成的，所以即使是使用了大量數(shù)學函數(shù)，主要的計算量也在FMA、MUL和ADD這種簡單指令。MUFU還是占比比較小。如果使用了內(nèi)置函數(shù)，則相對占比高一些，但多數(shù)時候還是需要相應(yīng)的簡單指令的配合。而且有些常用函數(shù)其實并沒有加速實現(xiàn)（比如sqrt，即使是intrinsic的__fsqrt_rn這種實現(xiàn)也非常復(fù)雜）。這里面肯定是有所考慮的。
MUFU的64bit版本（RCP64H和RSQ64H）其實只用了一個GPR輸入，估計覺得反正是近似，多后面32bit意義也不大。
FSETP
FSETP本身也沒有什么特殊的，就是根據(jù)一些浮點的比較操作設(shè)置一個bool變量。比如：
FSETP.GT.AND P1, PT, R1, R2, PT ;
FSETP.GEU.AND P1, PT, R1, RZ, PT ;
第一個是比較R1是否大于（greater than，GT）R2，然后順帶把結(jié)果與另一predicate做了與（AND）。GE是greater or equal，那GEU是什么呢？這其實是浮點比較中的一個特例。PTX把這個叫unordered floating-point comparisons，專門為了NAN的特殊形式。
如果熟悉IEEE 754，就知道NAN如果出現(xiàn)在比較浮點運算（a op b）中，那結(jié)果永遠是false，甚至x為nan時x==x都是false（甚至能用這個來判斷x是不是nan……）。所以PTX或SASS中有兩類浮點比較，一類是ordered，有nan就為false，一類是unordered，變成只要有nan就為true。如果兩個輸入都不是nan，那兩者就一樣。
這個邏輯有什么用呢？一些編譯器在處理比較運算的時候，會選擇一個canonical的形式（比如把a >= b統(tǒng)一改成 b<a，這樣就不用處理>=這種操作了），而對浮點操作這個變化是不正確的，就是因為有nan的存在。但是ordered a>=b 與 unordered b<a 兩者是等價的。
浮點的比較操作里還是有不少坑，和整數(shù)的比較還是差別很大。除了nan的不等于自己，還有一個就是+0和-0的二進制表示不一樣，但卻是相等的。denormal的情況我沒有仔細研究，要是會被flush的話，這里面也有很多需要注意的問題。
LDG/STG, LDS/STS, LDL/STL
內(nèi)存相關(guān)指令一般都比較復(fù)雜，特別是global操作，modifier特別多。這部分其實已經(jīng)超出指令集設(shè)計本身的范疇，更多的是consistency model的問題。這只討論指令輸入輸出中的一些點。
比如說LDG的有兩個使用UR做base的形式：
LDG.E.SYS R16, [R10.64+UR10+0x200] ;
LDG.E.CONSTANT.SYS R54, [R32.U32+UR4] ;
第一條中R10.64應(yīng)該表示它是個64bit的值，如果是32bit會用第二行的R32.U32這種形式。而后面的UR4，一般都是64bit。用32bit的好處就是一些簡單的地址運算可以用32bit的運算得到，比每次都做64bit操作要更劃算。而如果是warp內(nèi)uniform的base變化，可以直接改UR，UDP的計算與一般的ALU也是獨立的。這也算是一種減少計算強度的方式。
LDG/STG也有一些控制cache的邏輯，比如帶CONSTANT的modifier表示會在read-only L1中緩存。當然，這個具體操作其實每代架構(gòu)都會有所不同，但能用constant的話沒理由不用。
Shared Memory里有一個有意思的內(nèi)置left shift。比如下面的X4, X8, X16:
STS [R2.X4+0x400], R3 ;
STS.64 [R26.X8+UR4+0x10], R28 ;
LDS.U.128 R20, [R57.X16+UR4] ;
Shared memory因為地址窗口小，一般32bit地址足夠了。而且起始地址永遠是0，基址也常常是編譯期常數(shù)（靜態(tài)分配）。所以地址計算往往可以表示為immeBase + index * size。而元素大小常常是2的冪次，這樣就可以省去預(yù)先計算乘法或是移位，直接讓地址單元進行處理。這也算是減少計算強度的方法。
不過其實LDG/STG/LDL/STL應(yīng)該也可以用這個邏輯，特別是local memory，不知道為什么沒有。
BAR
BAR指令就是barrier，更確切的說是synchronize barrier，就是通常block內(nèi)用來同步的__syncthreads()函數(shù)。CUDA中叫barrier的術(shù)語挺多的，比如dependency barrier主要指DEPBAR指令，是用來等scoreboard的。memory barrier應(yīng)該是相當于memory fence，但是現(xiàn)在好像有專門的API操作。還有一種convergence barrier就比較底層了，是用來處理warp divergence的。
CUDA C一般只用一個barrier（就是BAR.SYNC 0x0后面的那個0x0），但實際上PTX或SASS里每個block最多可以用16個barrier，每次同步可以選擇同步到某個barrier。這樣同一個block的線程可以分塊同步到不同的barrier上去（BAR有warp計數(shù)的參數(shù)）。另外，BAR還支持arrive、sync模式，可以搭配出類似生產(chǎn)者-消費者的模型，具體可以參考PTX文檔。這個功能沒有C API，但是可以用inline PTX其實風險也不大。這其實也算是一個降低瓶頸資源壓力提高并行度的方法，與前面提到的用predicate做carry，還有跳轉(zhuǎn)用顯式GPR代替隱式棧，應(yīng)該算是一類思路。
CUDA 11之后有一個新功能叫Asynchronous Barrier。從實現(xiàn)上更偏軟件，但功能上沒看出它比BAR.arrive與BAR.sync的組合有什么更特別的地方。也許是scope更靈活？或是更適合CPU軟件移植和編譯器實現(xiàn)？
BAR不僅可以實現(xiàn)同步，還可以順帶計數(shù)或是做reduction，這個是可以從CUDA C調(diào)用的（參考）。從功能上講，用戶自己用shared memory實現(xiàn)應(yīng)該也可以，不過既然是順帶的，多少還是省點用戶的事。具體硬件實現(xiàn)上便不便宜，就不得而知了。
結(jié)語
指令集的設(shè)計其實是一個挺玄學的問題。好的指令集設(shè)計與好的產(chǎn)品之間，其實還隔著非常多其他因素。功能貼合市場需求，性能滿足要求，其實就具備了成為好產(chǎn)品的條件。而指令集設(shè)計本身依附于產(chǎn)品提供的功能和性能，它更多的是一個接口的角色：連接用戶與設(shè)備，讓功能和性能更容易發(fā)揮。但我覺得指令集并不是決定性因素。x86指令集現(xiàn)在常被吐糟是過時設(shè)計，即使它已經(jīng)不太強勢，但生命力仍然頑強。我覺得也不是光靠“兼容性”和“wintel壟斷”可以解釋得通的，至少它當前具有的功能和性能，一定是多年市場選擇和無數(shù)技術(shù)迭代的結(jié)果。當然，好的指令集設(shè)計，可以讓一個架構(gòu)更通用，性能可及性更好，硬件效率更高，從而更有生命力和競爭力。
二．CUDA微架構(gòu)與指令集-指令發(fā)射與warp調(diào)度
CUDA指令的發(fā)射和warp調(diào)度問題。
? 指令發(fā)射的基本邏輯，主要是指令發(fā)射需要滿足的條件，幾代架構(gòu)發(fā)射指令的一些簡單描述等等。
? control codes，主要是它與指令發(fā)射和warp調(diào)度的關(guān)系。
? warp Scheduler的功能。
? 峰值算力的計算方法和達成條件。
指令發(fā)射的基本邏輯
首先，簡單回顧一下CUDA程序的等級結(jié)構(gòu)。每個Kernel有一個grid，下面有若干個block，每個block有若干個warp。同一個block的warp只能在同一個SM上運行，但是同一SM可以可以容納來自不同block甚至不同grid的若干個warp。我們這里要聊的指令發(fā)射和warp調(diào)度的問題，就是指同一個SM內(nèi)同一個warp或是不同warp的指令之間是按照什么邏輯來調(diào)度運行的。
指令發(fā)射的一些基本邏輯：

1. 每個指令都需要有對應(yīng)的功能單元（Functional Unit）來執(zhí)行。比如執(zhí)行整數(shù)指令的單元，執(zhí)行浮點運算指令的浮點單元，執(zhí)行跳轉(zhuǎn)的分支單元等等。功能單元的個數(shù)決定了這種指令的極限發(fā)射帶寬（在沒有其他資源沖突時）。
2. 每個指令都要dispatch unit經(jīng)由dispatch port進行發(fā)射。不同的功能單元可能會共用dispatch port，這就意味著這些功能單元的指令需要通過競爭來獲得發(fā)射機會。不同的架構(gòu)dispatch port的數(shù)目和與功能單元分配情況會有一些差別。
3. 有些指令由于功能單元少，需要經(jīng)由同一個dispatch port發(fā)射多次，這樣dispatch port是一直占著的，期間也不能發(fā)射其他指令。比較典型的是F64指令和MUFU特殊函數(shù)指令。
4. 每個指令能否發(fā)射還要滿足相應(yīng)的依賴關(guān)系和資源需求。比如指令LDG.E R6, [R2] ;首先需要等待之前寫入R[2:3]的指令完成，其次需要當前memory IO的queue還有空位，否則指令也無法下發(fā)。還有一些指令可能有conflict的情況，比如shared memory的bank conflict，register的bank conflict，atomic的地址conflict，constant memory在同一warp內(nèi)地址不統(tǒng)一的conflict等等，這些都有可能導(dǎo)致指令re-issue（甚至cache miss也可能導(dǎo)致指令replay）。這些情況會不會重復(fù)占用dispatch port發(fā)射帶寬我暫時還沒仔細研究。
5. 在有多個warp滿足發(fā)射條件的情況下，由于資源有限，需要排隊等待發(fā)射，warp scheduler會根據(jù)一定的策略來選擇其中的一個warp進行指令發(fā)射。
6. 當前CUDA的所有架構(gòu)都沒有亂序執(zhí)行（Out of order），意味著每個warp的指令一定是按照運行順序來發(fā)射的。當然，有的架構(gòu)支持dual-issue，這樣可以有兩個連續(xù)的指令同時發(fā)射，前提是兩者不相互依賴，而且有相應(yīng)的空余資源（比如功能單元）供指令運行（對kepler來說不一定是不同類的功能單元，后面會具體分析）。另外一個顯而易見的要求是雙發(fā)射的第一個指令不能是分支或跳轉(zhuǎn)指令。
   接著我們要簡單講講Kepler，Maxwell，Turing三代架構(gòu)的指令發(fā)射和warp調(diào)度邏輯。先看看這三個微架構(gòu)典型chip的SM示意圖（圖來自各自whitepaper，注意SM的名稱有點變化）：
   

簡單對比一下其中的warp scheduler和dispatch unit：

每個warp scheduler每cycle可以選中一個warp，每個dispatch unit每cycle可以issue一個指令。幾個架構(gòu)的區(qū)別：

1. Kepler的SMX有192個core，但是core和warp沒有固定的從屬關(guān)系，相當于每個warp都可以運行在不同的32core組上（？這個需要進一步確認）。192=32*6，所以每個cycle最少需要發(fā)6個指令才能填滿所有core。可是4個warp scheduler只能選中4個warp。所以Kepler必須依賴dual issue才能填滿所有cuda core，而且由于core多且與warp沒有對應(yīng)關(guān)系，kepler的dual issue不一定是發(fā)給兩個不同的功能單元，兩個整數(shù)、兩個F32或是混合之類的搭配應(yīng)該也是可以的，關(guān)鍵是要有足夠多的空閑core。每個warp scheduler配了兩個dispatch unit，共8個，而發(fā)射帶寬填滿cuda core只要6個就夠了，多出來的2個可以用來雙發(fā)射一些load/store或是branch之類的指令。
2. Maxwell開始，SM的資源就做了明確的分區(qū)，每個warp都會從屬于某個分區(qū)，各分區(qū)之間有些功能單元（比如cuda core和F64單元，SFU）是不共享的。Maxwell的SMM有128個core，分成4組。每組有一個warp scheduler，2個dispatch unit，配32個CUDA core。這樣每cycle發(fā)一個ALU指令就可以填滿cuda core了，多出來的dispatch unit可以發(fā)射其他memory或是branch指令。由于功能單元做了分區(qū)，沒有冗余了，這樣Maxwell的dual-issue就不能發(fā)給同樣的功能單元了。
3. Turing的SM的core數(shù)減半，變成64個，但還是分成4個區(qū)，每個區(qū)16個core，配一個warp scheduler和一個dispatch unit。這樣兩個cycle發(fā)一個指令就足夠填滿所有core了，另一個cycle就可以用來發(fā)射別的指令。所以Turing就沒有dual-issue的必要了。從Volta開始，NV把整數(shù)和一些浮點數(shù)的pipe分開，使用不同的dispatch port。這樣，一些整數(shù)指令和浮點數(shù)指令就可以不用競爭發(fā)射機會了。一般整數(shù)指令用途很廣，即使是浮點運算為主的程序，也仍然需要整數(shù)指令進行一些地址和輔助運算。因此，把這兩者分開對性能還是很有一些幫助的。但是，這里還是要澄清一下，不是所有的浮點和整數(shù)都是分開的。這里貼一張Hotchips里NV介紹Turing的圖（圖上沒畫F64，應(yīng)該是和Tensor Core、F16在一塊）：

Volta和Turing的F64、F16和Tensor Core都是用的同一個dispatch port。然后F32一組（IMAD也放在這組，大概是要共享mantissa的那個乘法器），MUFU（就是特殊函數(shù)指令），其他ALU（包括整數(shù)算術(shù)運算和移位、位運算等等，但是不包括IMAD）一組，然后memory指令一組，branch指令一組，然后Turing的uniform datapath的指令一組，這些都各自有dispatch port。不同組的可以隔一個cycle發(fā)射，同組的就要看功能單元數(shù)目，最少是隔2個cycle。
Control codes
從Kepler開始，指令中就帶有Control codes。Kepler的架構(gòu)略顯久遠，格式又不太一樣，信息也不太全。Maxwell后control codes的功能更加豐富了，而且在Scott Gray的文章中非常具體詳盡的描述了各個域的含義和功能。Pascal的指令集與Maxwell基本一致，control codes也沒什么變化。Volta和Turing，包括最新的Ampere，只是把control codes編碼到每條指令中，具體內(nèi)容和bit對應(yīng)關(guān)系其實也沒有變，所以我這也不做架構(gòu)的區(qū)分。這里主要復(fù)述一下Scott Gray的表述，但主要側(cè)重它與指令發(fā)射和warp調(diào)度的關(guān)系。
這里我以Turing為例，選了一段程序，把對應(yīng)的control codes 寫在前面（這里我用的格式與Scott有點區(qū)別，主要是看起來更方便）：
1: [----:B------:R-:W0:-:S04] S2R R113, SR_CTAID.Y ;
2: [----:B------:R-:W1:-:S04] S2R R0, SR_CTAID.Z ;
3: [----:B------:R-:W3:-:S01] S2R R106, SR_TID.X ;
4: [----:B0-----:R-:W-:-:S02] IMAD.SHL.U32 R113, R113, 0x4, RZ ;
5: [----:B-1----:R-:W-:Y:S03] IMAD.SHL.U32 R0, R0, 0x4, RZ ;
6: [R—:B------:R-:W-:-:S02] IADD3 R107, R113.reuse, 0x1, RZ ;
7: [R—:B------:R-:W-:-:S01] IADD3 R109, R113.reuse, 0x2, RZ ;
8: [R-R-:B------:R-:W-:-:S01] IMAD R2, R113.reuse, c[0x0][0x1a8], R0.reuse ;
9: [----:B------:R-:W-:-:S01] IADD3 R111, R113, 0x3, RZ ;
我這里采用的顯示形式是類似R-R-:B------:R-:W-:-:S01這種，用冒號":“分隔開6個域：Register Reuse Cache（4bit，對應(yīng)4個slot，有reuse就寫R，沒有就”-"），Wait Dependency Barrier（6bit，B+6個數(shù)，有等待就寫上對應(yīng)的barrier號0-5，否則寫“-”），Read Dependency Barrier（3bit，R+設(shè)置的barrier號，不設(shè)置寫“-”），Write Dependency Barrier（3bit，W+設(shè)置的barrier號），Yield Hint Flag（1bit，“Y”表示yield，否則寫“-”），Stall Count（4bit，S+十進制的stall的cycle數(shù)）。
這里與scott的標記的區(qū)別一是顯式的寫了reuse，而不是只寫在后面匯編文本里。這樣看起來更直觀且容易看到slot的對應(yīng)關(guān)系，因為有的predicate夾在中間容易搞混。二是wait barrier的每個bit我拆開了，這樣肉眼更容易與前面設(shè)置barrier的R#或是W#對應(yīng)上。然后就是每個域前加了提示字符，BRWS之類，畢竟有的時候看久了也容易恍惚，這樣區(qū)分比較明顯，不容易搞混。
【更正1】：這里包括后文說的Dependency Barrier也許更合理的稱呼是Scoreboard，在profiler里有short scoreboard和long scoreboard之分，應(yīng)該指的就是這個。
【更正2】：scoreboard的編號我調(diào)整過一次。由于最開始我沿用了scott的1-6的編號方式，但后來我還是調(diào)整為更符合原語境的0-5。首先這是原編碼的值，其次DEPBAR后會顯式的接SB0~SB5這種數(shù)，所以后來我就統(tǒng)一改為0-5。
下面具體介紹control codes每個域的含義。
Register Reuse Cache
Register Reuse Cache有4bit。每個指令有4個slot，每個register的source operand的位置對應(yīng)一個slot（predicate好像不算）。我暫時還沒碰到4個source operand的指令，所以有一個bit好像一直沒用到。Reuse的用法：如果當前指令某個slot的register還會被下一個指令的同一個slot讀取，那就可以reuse當前指令讀取到的register內(nèi)容。Reuse的作用主要就是減少GPR的讀取，一來可以減少register bank conflict，二來應(yīng)該也能省一點功耗。比如前面代碼中Line6的IADD3的第一個源操作數(shù)是R113，而Line7的第一個源操作數(shù)也是R113，所以可以reuse。同理，line7、line8同樣位置的R113都可以reuse。但是line8的R0為什么reuse呢？我也沒搞懂。Reuse是唯一在官方反匯編中出現(xiàn)的control codes，但也有很多疑點。
我的幾個猜測：首先reuse是個cache，某種意義上是個hint，也就是說就算set了，reuse不了應(yīng)該也不至于出錯。第二，reuse cache的位置應(yīng)該是位于所謂的operand collector。這個應(yīng)該是很多功能單元公用的，所以不太可能是同種功能單元的指令才能reuse。另外，load store類的指令沒看到reuse，好像是不用collector。第三，如果切換到別的warp，register是不同的空間，那reuse cache就失效了。所以reuse應(yīng)該是需要當前warp的指令連續(xù)發(fā)射。第四，reuse既然是register的cache，那原來register的位置出現(xiàn)了constant memory或是immediate會不會使reuse失效呢？感覺好像不會。那RZ會不會呢？UR會不會呢？我也沒仔細研究過。
Reuse這里還有很多奇怪的問題，有的我懷疑是編譯器的bug。比如這種：
1: [R—:B------:R-:W-:-:S02] FSETP.GT.AND P1, PT, R9.reuse, c[0x0][0x18c], PT ;
2: [----:B------:R-:W-:-:S02] LEA.HI.X.SX32 R4, R4, c[0x0][0x164], 0x1, P3 ;
3: [R—:B------:R-:W-:-:S02] LEA R5, P3, R0.reuse, R5, 0x2 ;
4: [----:B------:R-:W-:-:S02] FSEL R9, R9, c[0x0][0x18c], !P1 ;
5: [----:B------:R-:W-:-:S02] LEA.HI.X R4, R0, R4, R3, 0x2, P3 ;
其中Line1的R9只能在Line4中得到reuse，但之前的這個slot已經(jīng)被Line2的R4用過了，應(yīng)該是沒法reuse的。Line3也是一樣，R0只能在Line5中被reuse，但中間被Line4打斷了。如果我們把Line4提到第2行前面去，這好像就順理成章了：
1: [R—:B------:R-:W-:-:S02] FSETP.GT.AND P1, PT, R9.reuse, c[0x0][0x18c], PT ;
4: [----:B------:R-:W-:-:S02] FSEL R9, R9, c[0x0][0x18c], !P1 ; // 原先在第4行，依賴Line1的輸出P1，如果放在這就要等barrier
2: [----:B------:R-:W-:-:S02] LEA.HI.X.SX32 R4, R4, c[0x0][0x164], 0x1, P3 ;
3: [R—:B------:R-:W-:-:S02] LEA R5, P3, R0.reuse, R5, 0x2 ;
5: [----:B------:R-:W-:-:S02] LEA.HI.X R4, R0, R4, R3, 0x2, P3 ;
由于Line4的FSEL用到了Line1的輸出P1，為了把依賴指令移到后面以減少stall，編譯器做了相應(yīng)的指令調(diào)度，但是reuse卻沒有重新生成。所以我懷疑這個是編譯器的bug，存疑中。
Wait Dependency Barrier
Wait Dependency Barrier有6bit，每個bit表示是否需要等待對應(yīng)的dependency barrier。每個線程有6個dependency barrier，每個barrier都可以被后面的Read或Write操作設(shè)置上。設(shè)置wait dependency barrier是等待依賴的其中一種方式。SASS里面還有一個對應(yīng)的指令，如DEPBAR.LE SB0, 0x0, {2,1} ;，兩者自由度不太一樣。control codes里設(shè)置的barrier只能是bool形式，要么dependency resolved，要么就是not ready。而DEPBAR可以等待有計數(shù)的barrier。比如發(fā)了8個memory指令，其實設(shè)置的是同一個dependency barrier，每發(fā)出一個計數(shù)加1，每回來一個計數(shù)減1。這樣DEPBAR可以等計數(shù)降到6就說明前面2個指令已經(jīng)到位了，對應(yīng)的load結(jié)果就能用了。而如果用control codes來等待，就只能等計數(shù)降到0，也就是8個都ready才行。
dependency barrier有一個和分支指令強相關(guān)的地方，比如不確定的跳轉(zhuǎn)指令（帶predicate的BRA，或是BRX這種指令）需要等待當前所有已設(shè)置的dependency barrier都到齊才行。否則后面多個分支的代碼可能用到這個barrier，但又不一定都會等待。這個可能是編譯器在處理上有一些圖方便的地方。有些情況是可以把wait后移到對應(yīng)使用指令上的，這樣延遲更容易被隱藏。只是有時候編譯器拿不到足夠的信息，為保證正確性就統(tǒng)一在跳轉(zhuǎn)的時候等了。
Read Dependency Barrier
Read dependency barrier有3bit，表示需要設(shè)置的6個barrier中對應(yīng)的索引（0~5，對應(yīng)barrier 1-6，如果不需要設(shè)置barrier，就設(shè)置為0b111）。Read dependency barrier主要是一些指令不會在一開始就把所有操作數(shù)讀進去，所以需要hold住GPR的值，防止后面的指令在它讀取其內(nèi)容之前把GPR改掉。使用Read dependency barrier的主要是memory類的指令，但是一些轉(zhuǎn)換指令如F2I/I2F之類好像偶爾也能見到。
Write Dependency Barrier
Write dependency barrier與read dependency很類似，也是3bit，后面跟barrier索引。注意Read和Write兩者用的dependency barrier資源是一樣的，也都是上面wait的那6個。Write dependency barrier比較好理解，就是某個指令要把操作結(jié)果保存到某個GPR或是predicate中，使用barrier進行保序可以防止出現(xiàn)data race。不過這主要針對的是不定長latency指令。如果一個指令的latency是確定的（或者有不太長的上限），那用后面提到的stall cycle停足夠長時間就可以保證沒有race。
Yield Hint Flag
Yield hint flag是1bit。如果Yield，就表示下一個cycle會優(yōu)先發(fā)射其他warp的指令。前面聊reuse的時候我們也說了，reuse是需要連續(xù)發(fā)射同一個warp的指令的。所以reuse和yield是不會聯(lián)用的。另外，SASS有一個專門的指令YIELD，感覺上是一樣的功能（【注】：經(jīng)評論區(qū)提醒，功能還是不一樣的，具體請移步評論區(qū)）。Yield的存在仿佛暗示warp scheduler不是round robin的選擇warp，而是傾向于一直往同樣的warp里發(fā)射指令（如果可以一直發(fā)射的話，有stall肯定就盡量切走了）。但是這個東西我也不太確定，沒仔細測過。如果真是這樣，yield的作用就是保持各個warp之間的進度均衡，否則在barrier之類的指令上會有較大的等待開銷。而且如果退出時間差很大，也會導(dǎo)致一些資源不能盡快回收以容納新的block。
【補充】：根據(jù)評論區(qū)所說，也許yield這個bit就是stall count的高位。只是假如這個bit不為0，那stall的cycle會>16，相當于warp被切換的概率也會大大增加。兩者之間的具體含義應(yīng)該還和具體指令有關(guān)，不同類型的指令也許是不一樣的。這個問題我還沒有具體研究過。先存疑。
Stall Count
Stall count有4bit，表示當前指令后需要stall指令發(fā)射的cycle數(shù)，然后再決定是不是要繼續(xù)發(fā)射。這個cycle數(shù)受到極限發(fā)射帶寬的約束，很多時候可以用來反推功能單元的分組和數(shù)目。比如Maxwell下有雙發(fā)射，所以可以stall 0 cycle，在反匯編中會用大括號組合起來，比如這種：
/0008/ MOV R1, c[0x0][0x20] ; /* 0x4c98078000870001 /
/0010/ { MOV R5, c[0x0][0x148] ; / 0x4c98078005270005 /
/0018/ S2R R0, SR_TID.X }
/ 0xf0c8000002170000 */
Kepler的雙發(fā)射不限定功能單元，Turing架構(gòu)沒有雙發(fā)射，就都沒有這種形式了。
我們再找個Turing的例子來看一下：
[----:B------:R-:W-:-:S01] IADD3 R43, -R43, R28, R36 ; // STS與IADD3不同功能單元，stall 1 cycle
[----:B–2—:R0:W-:-:S04] @!P3 STS.128 [R31.X16], R8 ; // 兩個STS間至少stall 4 cycle，
[----:B—3–:R1:W-:-:S04] @!P5 STS.128 [R31.X16+0x2100], R16 ;
[----:B----4-:R2:W-:-:S01] @!P4 STS.128 [R51.X16+0x2200], R12 ;
[R-R-:B0-----:R-:W-:-:S01] IADD3 R11, -R28.reuse, 0x20, R43.reuse ; // IADD3是Integer pipe，IMAD是F32 pipe，不共dispatch port，stall 1 cycle
[----:B------:R-:W-:-:S01] IMAD.IADD R8, R43, 0x1, -R28 ;
[----:B------:R-:W-:Y:S02] IADD3 R9, -R28, 0x10, R43 ; // IADD3和ISETP同屬普通Integer pipe，stall 2 cycle
[R—:B------:R-:W-:-:S01] ISETP.GE.U32.AND P3, PT, R0.reuse, R11, PT ; // ISETP與LDG屬不同功能單元，stall 1 cycle
[----:B-1----:R0:W5:-:S01] @!P2 LDG.E.128.CONSTANT.SYS R16, [R24+0x200] ;
前面我們已經(jīng)講過，Turing的兩個周期發(fā)一個ALU指令就可以用滿cuda core了，所以同組ALU指令至少stall 2 cycle。如果兩個指令間只stall 1 cycle，說明這兩個指令應(yīng)該分屬不同的功能單元（或者說不共用dispatch port），可以分開發(fā)射。如果stall的時間更長，說明其發(fā)射帶寬比較低，比如LDG/LDS/STG/STS這種memory指令都是4 cycle才能發(fā)一個。但是也有一些奇怪的例子：
[----:B------:R-:W-:-:S01] IMAD.WIDE.U32 R2, R5, 0x3, R2 ;
[----:B–2—:R-:W0:-:S01] I2F.S16 R4, R4 ;
[----:B—3–:R-:W1:-:S01] I2F.S16 R6, R6 ;
[----:B----4-:R-:W2:-:S01] I2F.S16 R7, R7 ;
[----:B0-----:R-:W-:-:S01] FADD R8, R4, -c[0x0][0x194] ;
[----:B------:R-:W-:-:S01] SHF.R.S32.HI R4, RZ, 0x1f, R5 ;
[----:B-1----:R-:W-:-:S02] FADD R9, R6, -c[0x0][0x198] ;
I2F居然是挺機關(guān)槍，每cycle都可以發(fā)。可是按programming guide中的Instruction Throughput表，I2F如果不涉及F64，是1/4的throughput。難道是有特別的queue做buffer？這個也沒太明白，有待進一步研究。
其他
Control codes一個比較容易忽視的問題是它與predicate是獨立的。也就是說不管加不加predicate，control codes的作用是不會改變的。因為本身control codes很多東西是編譯期決定的，如果按運行期的predicate來定是否啟用control codes，有些代碼的正確性就容易出問題。
Control codes是個很復(fù)雜的問題，有些東西真是純逆向了。我也沒仔細研究過。大家有興趣可以看看scott gray的maxas的實現(xiàn)，或許會有一些啟發(fā)。
Warp Scheduler
Warp scheduler的作用就是管理一系列的warp，在那些滿足條件的warp中選中一個來發(fā)射指令。就緒可以發(fā)射指令的warp就是eligible，不滿足發(fā)射條件的warp就是stalled。導(dǎo)致warp不能發(fā)射指令的原因有很多種。我根據(jù)NSight Visual Studio Edition中Issue Stall Reasons中的描述，大致搬運翻譯一下：
? Pipeline Busy：指令運行所需的功能單元正忙。
? Texture: Texture單元正忙，或者說已經(jīng)下發(fā)的request過多。
? Constant：Constant cache的miss。一般說來，多數(shù)情況下constant cache的hit rate還是很高的，所以一般只會在第一次access的時候miss。
? Instruction Fetch：Instruction cache的miss。與constant cache的miss類似，一般只有第一次運行到的地方才容易miss。比如BRA新跳轉(zhuǎn)到的地方，或是Instruction cache的cache line的邊界處。
? Memory Throttle：有大量memory操作尚未完成，導(dǎo)致memory指令無法下發(fā)。可以通過合并memory transactions來緩解.
? Memory Dependency: 由于請求資源不可用或是滿載導(dǎo)致load/store無法執(zhí)行，可以通過內(nèi)存訪問對齊和改變access pattern來緩解。這個與memory throttle的細微差別我還沒仔細研究過。
? Synchronization：warp在等待同步指令，如cuda C里的_syncthreads()，對應(yīng)SASS里的BAR指令。
? Execution Dependency：輸入依賴關(guān)系沒解決。簡單說，就是輸入值未就緒，就是在等control codes里的dependency barrier。單個warp內(nèi)通過增加ILP可以減少依賴型stall。如果ILP不夠用，這個stall就會形成額外的latency，只能用TLP來隱藏了。
在Profiler里提供的Turing的performance counter里，還有兩個當前warp不能發(fā)射指令的原因:
? stall_not_selected: warp當前雖然eligible，eligible的warp超過一個，當前的未被選中，所以不能發(fā)射。
? stall_sleeping: 這個一般是用戶自己調(diào)用sleep功能讓該warp處于睡眠狀態(tài)。
貼一個Nsight里的stall reason統(tǒng)計圖：

Warp Scheduler的另一個關(guān)鍵功能是在eligible里選一個來發(fā)射。很多早期的書上都說選warp是round-robin，就是所謂的輪詢，發(fā)一個換一個。當前的幾代架構(gòu)，至少maxwell之后，我覺得應(yīng)該不是這個策略了。前面聊reuse，yield和stall的時候也提到了，如果是round-robin，這些東西都會顯得很奇怪了。所以我感覺它應(yīng)該是比較aggresive的往同一個warp發(fā)射指令，除非stall了。當然，中間如果沒有yield，那stall 2個cycle的時候中間那個cycle能去發(fā)射別的warp嗎？這個我也有點迷，有機會再仔細研究下。
【補充】：我仔細再想了想，stall 2cycle的時候中間那個cycle應(yīng)該是可以發(fā)射其他不用operand collector的單元，比如memory和branch。網(wǎng)上看到一些說法，有些S2R的指令好像也是走的memory的pipe，那應(yīng)該也可以發(fā)射。這個應(yīng)該是可以寫一個micro-benchmarking的case驗證一下，只是我手上沒有Turing的卡，也沒法測了。但是我感覺只要它不影響reuse的cache，那就應(yīng)該可以抽空發(fā)射。這也是一個比較符合性能需求的模式。
Eligible的warp數(shù)是影響峰值性能的關(guān)鍵表征之一，如果每個cycle都至少有一個eligible warp，那功能單元基本就會處于滿載狀態(tài)，性能一般也會比較好。這也是occupancy真正起作用的方式。

關(guān)于峰值算力的問題
最后要略微展開聊一下峰值算力的問題。NV的GPU經(jīng)常用CUDA Core這個詞來表示算力強弱，這其實是Int32/Float32功能單元的marketing叫法。算力評估常用的單位是FLOPS，表示FLoat OPerations per Second（Flops有時候也用作Flop的復(fù)數(shù)，注意鑒別）。對于F32而言，FADD/FMUL都是一個指令一個flop，FFMA一個指令同時算乘加所以是兩個flop。所以一般NV的GPU的F32峰值算力計算方法為：
SM數(shù) * 每SM的Core數(shù) * 2 * 運行頻率
最后結(jié)果常用GFlops或是TFlops表示。其中乘以“2”是因為FFMA是兩個flop。比如說Maxwell架構(gòu)每個SM有128個CUDA core，每個SM每cycle可以發(fā)射128條Int32或Float32指令（兩種指令不能同時發(fā)射，所以是或）。Maxwell架構(gòu)的GTX 980有16個SM，共16 * 128=2048個CUDA core，每個cycle能做2048 * 2Flops/FFMA=4096 Flops。980的base clock是1.126GHz，相當于每個cycle是1/1.125G秒，每個Cuda Core每秒可以發(fā)射1.126G條指令，整個GPU就是1.126G*4096=4.5TFlops。放在一起算，就是峰值算力等于：
16 SM * 128 Core/SM * 2 Flop/Core/Cycle * 1.126G Cycle/second = 4.5TFlops.
當然，商家為了宣傳，常用boost clock算峰值算力，非公版的頻率也會有些差別，所以這個值會有些變化。另外，這里用的是F32浮點峰值做例子，如果你的任務(wù)不需要浮點運算或是精度不是F32，這個值就意義不大，需要轉(zhuǎn)換成你需要的那個操作。現(xiàn)在AI處理器常常宣傳峰值是多少FLOPs，或是多少IOPs，一般也會限定是F32，F16或是I8之類。因為每種操作對應(yīng)的指令是不一樣的，峰值當然也可能不一樣。頂級HPC計算卡F64一般是F32的一半，但消費級顯卡F64多數(shù)會有閹割。如果沒有TensorCore而用packed F16（把兩個F16塞到一個32bit GPR里同時運算），F16峰值性能通常是F32的兩倍。有TensorCore時則要另行計算，要看具體TensorCore的數(shù)目和指令帶寬，還有能不能和其他指令同時發(fā)射等。其實不看Tensor core的話，滿血版一般有：F64:F32:F16=1:2:4，正好與占用的GPR成反比，這個其實是與GPR的帶寬有很大的關(guān)聯(lián)的，一般滿血版的卡的功能單元配比就會盡量按極限的GPR帶寬來設(shè)計。
這里貼一個Ampere的white paper中V100和A100的幾種峰值性能對比：
要達到F32的峰值性能，需要滿載發(fā)射FFMA指令，這是很苛刻的條件。首先，其他與FFMA共用dispatch port的指令，每發(fā)射一個都會擠占FFMA的發(fā)射機會。其次，由于多數(shù)情況下數(shù)據(jù)要從memory中來，而memory操作比ALU慢很多，常常導(dǎo)致指令操作數(shù)無法就緒，從而有些周期沒有FFMA指令可發(fā)。同時還有其他一些overhead或是occupancy問題導(dǎo)致有些SM無法滿載，從而無法達到峰值。一般說來，實際應(yīng)用中，較大尺寸的矩陣乘法（GEMM）是難得的能接近峰值性能的程序，有些實現(xiàn)能到98%峰值的效率。但多數(shù)實際應(yīng)用效率都遠不及此，很多memory bound程序能到10%就很不錯了。超算TOP500排名中，多數(shù)HPL效率都是50_{60%左右，更接近實際應(yīng)用的HPCG效率一般都在2}3%左右。雖有規(guī)模大導(dǎo)致的互聯(lián)開銷原因，但總體來講實際應(yīng)用的峰值性能離極限值還是差距很大的。
每種類型的指令都有一個峰值性能，那是不是能同時達到呢？基本是不能。對于共用dispatch port就不說了，要相互競爭發(fā)射機會，發(fā)射一條這種指令就少發(fā)射一條那種指令，所以顯然不能同時到達峰值。如果是不同的dispatch port呢？理論上可以，但是實際上也會比較難。比如說Turing的F32和I32，首先I32的2IOP指令I(lǐng)MAD是和F32一伙的，相互競爭dispatch port，所以兩個不能同時到達峰值。剩下的IADD3或是LEA之類的指令理論上可以與F32的并行，倒是有機會沖一沖。只不過多數(shù)實際應(yīng)用中很難做到這么好的運算配比，而且register的bank conflict之類應(yīng)該也會大大限制這兩種指令的同時運行。另外，即使兩者配合完美，它還是需要省出一些發(fā)射帶寬給其他配套指令（比如memory load），不可能完全占滿。
本文將兩篇文章整理了一下。
參考文章鏈接如下：
https://zhuanlan.zhihu.com/p/391238629
https://zhuanlan.zhihu.com/p/166180054

總結(jié)

以上是生活随笔為你收集整理的GPU指令集技术分析的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。