GPU指令集技术分析
GPU指令集技術分析
本文將兩篇文章整理了一下。
參考文章鏈接如下：
https://zhuanlan.zhihu.com/p/391238629
https://zhuanlan.zhihu.com/p/166180054
一．GPGPU- 指令執行設計
本節主要內容：
? GPGPU指令執行簡介
o GPGPU指令執行流水線
o GPGPU指令執行吞吐的影響因素
o GPGPU指令執行的特點
? 指令設計中的一些原則與思路
o 指令長度的問題
o 指令集設計與ILP的一些相關性
o 復合操作與附加操作
o 立即數操作數和Constant Memory操作數
o 格式的通用性與信息具體化
? 簡單聊聊一些具體的指令
o FMA, MUL, ADD 系列
o IMAD, LEA, IADD3
o LOP3
o MUFU
o FSETP
o LDG/STG, LDS/STS, LDL/STL
o BAR
? 結語
二．CUDA微架構與指令集-指令發射與warp調度
CUDA指令的發射和warp調度問題。
? 指令發射的基本邏輯，主要是指令發射需要滿足的條件，幾代架構發射指令的一些簡單描述等等。
? control codes，主要是它與指令發射和warp調度的關系。
? warp Scheduler的功能。
? 峰值算力的計算方法和達成條件。
現在分開來敘述
一．GPGPU- 指令執行設計
GPGPU指令吞吐和指令集設計的一些問題。NV GPU的機器碼指令集叫SASS。
指令集是微架構與用戶對接的途徑。指令集相當于硬件提供給軟件的API（或者也可以認為指令集是輸入前端，微架構是執行后端）。按照
GPGPU指令執行簡介
GPGPU指令執行流水線
首先，先簡單介紹一下通用處理器的指令執行邏輯。對RISC V一種簡單的五級流水線實現的描述：
- Instruction fetch cycle (IF)：主要是獲得Program Counter對應的指令內容。
- Instruction decode/register fetch cycle (ID)：解碼指令，同時讀取輸入寄存器的內容。由于RISC類指令的GPR編碼位置相對固定，所以可以直接在解碼時去讀取GPR的值。
- Execution/effective address cycle (EX)：這是指令進入執行單元執行的過程。比如計算memory地址的具體位置（把base和offset加起來），執行輸入輸出都是GPR或輸入含立即數的ALU指令，或者是確認條件跳轉指令的條件是否為真等。
- Memory access (MEM)：對于load指令，讀取相應的內存內容。對于store指令，將相應GPR的值寫入到內存地址處。RISC類指令集通常都是load-store machine，ALU指令不能直接用內存地址做操作數（只能用GPR或立即數），因而通常ALU指令沒有memory access。
- Write-back cycle (WB)：將相應的ALU計算結果或memory load結果寫入到相應的GPR中。
與復雜的CISC指令集相比，多數GPGPU的指令集還是比較接近load-store machine，總體來說與RISC更相似一些。GPGPU典型微架構可以簡單表示為下圖：
GPGPU指令執行流程示意圖
這個圖是微架構的模塊示意圖，而非流水線示意圖。實際執行流程中Fetch、Decode、Execution這三步是必須的，而Memory access顯然只針對memory指令，write-back則只針對需要寫回的指令（比如memory load，帶GPR輸出的指令等）。
注：流水線的配置與ALU的latency有很大的關系。比如Volta前FFMA的延遲是6cycle，Volta及之后FFMA的延遲是4cycle，這絕對與流水線的改進有關。不過，這里的Latency并不是所有流水線的級數。因為Latency在程序中的表現形式是：一個指令發射后，其結果需要多少周期才能就緒（也就是能被其他指令使用）。兩個back-to-back dependent的ALU指令（比如FFMA R0, R1, R2, R0; FFMA R0, R3, R4, R0;），前一個FFMA只要在第二個FFMA讀取操作數之前把結果寫回GRF，那后一個FFMA就可以得到正確值。對應到上面的5級流水線形式，就是前一個指令的WB要在后一個指令的ID前執行完就行（相當于4cycle延遲），最開始的IF那一級是不影響的。CPU對于這種形式的依賴還有更激進的旁路邏輯（forwarding），可以直接在前一個ALU的EX后把結果直接送給后一個ALU的EX當輸入，從而減少流水線的bubble，提高性能。NV的GPU應該是沒有這么緊湊的forwarding，但是NV的operand collector可以作為一個公共的操作數中轉站，理論上前一個ALU的結果寫回到operand collector就可以被下一個ALU看到了，不一定要回到GRF。
由于GPU運行模型的復雜性，在Decode后Execution前，還有大量其它的步驟：比如scoreboard的判斷（主要用來保證每個指令的執行次序，防hazard），warp Scheduler的仲裁（多個eligible warp中選一個），dispatch unit和dispatch port的仲裁（發射或執行資源沖突時需要等待），還可能有讀取作為立即數操作數的constant memory，讀取predicate的值生成執行mask，等等。在執行過程中，也有很多中間步驟，比如輸入GPR操作數的bank仲裁，跳轉指令要根據跳轉目標自動判斷divergence并生成對應的mask，訪存指令要根據地址分布做相應的請求廣播、合并等等。在寫回這一級的時候，由于一些指令的完成是異步的（比如一些內存指令），所以也可能需要GPR端口的仲裁，等等。
當然，步驟雖然多而瑣碎，但未必都會新增單獨的一級流水。GPU應該是為了簡化設計和節約功耗，不愿意把流水線拉得太細長，因而很多操作都是塞在同一級流水線里，各種組合邏輯非常復雜。這樣也導致它的主頻往往就不能太高。比如最近幾代NV GPU旗艦和次旗艦的主頻：
可以看到主頻基本都在1~2 GHz之間，次旗艦的頻率往往比旗艦要稍高一些（這里選的公版頻率，但非公也大致是這個趨勢），有些低端芯片頻率可能還會更高一點。而如今（2021年）常見的桌面端x86 CPU，基準頻率3~4 GHz，最大睿頻4~5 GHz是很尋常的事。很多Arm CPU的大核，主頻也能接近甚至超過3 GHz。當然，這么比也許不是特別公平。因為多數獨立GPU的功耗很大，會極大的限制頻率提升。一些眾核CPU的頻率也會比少核版的降一些，不過差別不會太大（類似上面GPU的旗艦與次旗艦的關系）。但即使算上這些，GPU的主頻比常見CPU的主頻還是顯著低一些（實際上帶核芯顯卡的CPU里，GPU頻率往往也是顯著小于CPU頻率）。這里面具體的因果關系我也不是特別明白，感覺肯定還是有些故事的~
GPGPU指令執行吞吐的影響因素
指令執行吞吐一般指的是每個時鐘周期內可以執行的指令數目，不同指令的吞吐會有所不同。通常GPU的指令吞吐用每個SM每周期可以執行多少指令來計量。對于多數算術邏輯指令而言，指令執行吞吐只與SM內的單元有關，整個GPU的吞吐就是每個SM的吞吐乘以SM的數目。而GPU的FMA指令（通常以F32計算）往往具有最高的指令吞吐，其他指令吞吐可能與FMA吞吐一樣，或是只有一半、四分之一等等。所以很多英文文檔會說FMA這種是full throughput，一半吞吐的是half rate，四分之一的是quarter rate等。當然，有些微架構下也會有1/3、1/6之類非2的冪次的比率。NV GPU近幾代微架構的常見指令吞吐如下：
CUDA算術邏輯指令吞吐表
從圖中可以發現，指令吞吐不僅與指令類型有關，還與微架構具體設計實現有關。主要會受到以下一些因素的影響：
- 功能單元的數目。絕大多數指令的功能都需要專用或共享的硬件資源去實現，設計上配置的功能單元多，指令執行的吞吐才可能大。顯然，只有最常用的那些指令，才能得到最充分的硬件資源。而為了節約面積，很多指令的功能單元會相互共享，所以他們的吞吐往往也會趨于一致。比如浮點的FFMA、FMUL都要用到一個至少24bit的整數乘法器（32bit浮點數有23bit尾數，小數點前還有1bit）。以前一些處理器有24bit的整數乘法指令，兩者乘法器就可以共用，從而具有相同的吞吐（不過NV最近幾代好像都沒有這個指令，ptx以及內置函數的24bit乘法應該是多個指令模擬的）。而FADD雖然用不上那個乘法器，但可以與FFMA共用那個很寬的加法器，以及一些通用的浮點操作（特殊數的處理，subnormal flush之類）。32bit的整數乘法因為需要更寬的乘法器，有的就不會做成full throughput，甚至可能被拆分成多個指令（比如Maxwell和Pascal用三個16bit乘法指令XMAD完成一次32bit整數乘法）。Turing的IMAD應該是有意識的加寬了，所以32bit的IMAD與FFMA吞吐一樣，但印象中帶64bit加數的IMAD應該還是一半。再比如一些超越函數指令（MUFU類，比如rcp，rsq，sin，exp之類），由于實際使用量相對不會太頻繁，多數是1/4的throughput。
- 指令Dispatch Port和Dispatch Unit的吞吐。這個在之前的專欄文章也詳細講過。一個warp的指令要發射，首先要eligible，也就是不要因為各種原因stall，比如指令cache miss，constant immediate的miss，scoreboard未就位，主動設置了stall count等等。其次要被warp scheduler選中，由Dispatch Unit發送到相應的Dispatch Port上去。Kepler、Maxwell和Pascal是一個Warp Scheduler有兩個Dispatch Unit，所以每cycle最多可以發射兩個指令，也就是雙發射。而Turing、Ampere每個Warp Scheduler只有一個Dispatch Unit，沒有雙發射，那每個周期就最多只能發一個指令。但是Kepler、Maxwell和Pascal都是一個Scheduler帶32個單元（這里指full-throughput的單元），每周期都可以發新的warp。而Turing、Ampere是一個Scheduler帶16個單元，每個指令要發兩cycle，從而空出另一個cycle給別的指令用。最后要求Dispatch Port或其他資源不被占用，port被占的原因可能是前一個指令的執行吞吐小于發射吞吐，導致要Dispatch多次，比如Turing的兩個FFMA至少要stall 2cycle，LDG之類的指令至少是4cycle。更詳細的介紹大家可以參考之前的專欄文章。
- GPR讀寫吞吐。絕大部分的指令都要涉及GPR的讀寫，由于Register File每個bank每個cycle的吞吐是有限的（一般是32bit），如果一個指令讀取的GPR過多或是GPR之間有bank conflict，都會導致指令吞吐受影響。GPR的吞吐設計是影響指令發射的重要原因之一，有的時候甚至占主導地位，功能單元的數目配置會根據它和指令集功能的設計來定。比如NV常用的配置是4個Bank，每個bank每個周期可以輸出一個32bit的GPR。這樣FFMA這種指令就是3輸入1輸出，在沒有bank conflict的時候可以一個cycle讀完。其他如DFMA、HFMA2指令也會根據實際的輸入輸出需求，進行功能單元的配置。
- 很多指令有replay的邏輯，這就意味著有的指令一次發射可能不夠。這并不是之前提過的由于功能單元少，而連續占用多輪dispath port，而是指令處理的邏輯上有需要分批或是多次處理的部分。比如constant memory做立即數時的cache miss，memory load時的地址分散，shared memory的bank conflict，atomic的地址conflict，甚至是普通的cache miss或是TLB的miss之類。根據上面Greg的介紹，Maxwell之前，這些replay都是在warp scheduler里做的，maxwell開始將它們下放到了各級功能單元，從而節約最上層的發射吞吐。不過，只要有replay，相應dispath port的占用應該是必然的，這樣同類指令的總發射和執行吞吐自然也就會受影響。
幾個需要注意的點： - 指令發射吞吐和執行吞吐有時會有所區別。有些指令有專門的Queue做相應的緩存，這樣指令發射的吞吐會大于執行的吞吐。這類指令通常需要訪問競爭性資源，比較典型的是各種訪存指令。但也有一些ALU指令，比如我們之前提過的Turing的I2F只有1/4的吞吐，但是可以每cycle連發（也就是只stall 1cycle）。不過多數ALU指令的發射吞吐和執行吞吐是匹配的。
- 要注意區分指令吞吐與常說的FLOPS或是IOPS的區別。通常的FLOPS和IOPS是按乘法和加法操作次數計算，這樣FMUL、FADD是一個FLOP，FFMA是兩個FLOP。這也是通常計算峰值FLOPS時乘2的由來。但是有些指令，可以計算更多FLOP。比如HFMA2 R0, R1, R2, R3;可以同時算兩組F16的FMA，相當于每個GPR上下兩個16bit分開獨立計算（類似于CPU的SIMD指令），所以SM86以前的架構HFMA2的指令吞吐與FFMA是一樣的，只是每條指令算4個F16的FLOP，而FFMA是2個F32的FLOP。這也就是TensorCore出現前F16的峰值通常是F32兩倍的原因。DFMA由于輸入寬度比FFMA再翻倍，所以功能單元做成一半就能把GPR吞吐用滿（這里說的是滿配Tesla卡，消費卡F64常有縮減）。因此，在TensorCore出現以前，通常的Tesla卡HFMA2、FFMA的指令吞吐一樣，DFMA吞吐是一半，而看峰值FLOP就是H:F:D=4:2:1的關系。TensorCore出現后，指令（比如HMMA）本身的吞吐和指令入口的GPR輸入量沒有變化，但由于同一個warp的指令可以相互共享操作數數據，一個指令能算的FLOP更多了，因而峰值又提高了。當然，這里說的是一般情況，實際上根據產品市場定位的不同，有些功能可能會有所調整。
- SM86（Ampere的GTX 30系列）的F32比較另類。Turing把普通ALU和FFMA（包括FFMA、FMUL、FADD、IMAD等）的PIPE分開，從而一般ALU指令可以與FFMA從不同的Dispatch Port發射，客觀上是增加了指令并行度。NVIDIA對CUDA Core的定義是F32核心的個數，所以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指令執行的特點
與傳統的x86 CPU相比，GPGPU在指令執行的邏輯上有很多獨特的地方。
靜態資源分配：GPU有一個很重要的設計邏輯是盡量減少硬件需要動態判斷的部分。GPU的每個線程和block運行所需的資源盡量在編譯期就確定好，在每個block運行開始前就分配完成（Block是GPU進行運行資源分配的單元，也是計算Occupancy的基礎）。典型的運行資源有GPR和shared memory。GPU程序運行過程中，一般也不會申請和釋放內存（當然，現在有device runtime可以在kernel內malloc和free，供Dynamic Parallelism用，但這個不影響當前kernel能用的資源）。CPU在運行過程中有很多所需的資源是動態調度的。比如，x86由于繼承了祖上編碼的限制，ISA的GPR數目往往比物理GPR少，導致常常出現資源沖突造成假依賴。實際運行過程中，通常會有register renaming將這些ISA GPR映射到不同的物理GPR，從而減少依賴（有興趣的同學可以研究下tomasulo算法）。GPU沒有這種動態映射邏輯，每個線程的GPR將一一映射到物理GPR。由于每個線程能用的GPR通常較多，加上編譯器的指令調度優化，這種假依賴對性能的影響通常可以降到很低的程度。
每個block在運行前還會分配相應的shared memory，這也是靜態的。這里需要明確的是，每個block的shared memory包括兩部分，寫kernel時固定長度的靜態shared memory，以及啟動kernel時才指定大小的動態shared memory。雖然這里也分動靜態，但指的是編譯期是否確定大小，在運行時總大小在kernel啟動時已經確定了，kernel運行過程中是不能改變的。
其實block還有一些靜態資源，比如用來做block同步的barrier，每個block最多可以有16個。我暫時沒測試到barrier的數目對Occupancy的影響，也許每個block都可以用16個。另一種是Turing后才出現的warp內的標量寄存器Uniform Register，每個warp 63個+恒零的URZ。因為每個warp都可以分配到足額，應該對Occupancy也沒有影響。另外每個線程有7個predicate，每個warp有7個Uniform predicate，這些也是足額，也不影響Occupancy。
GPU里還有一種半靜態的stack資源，通常也可以認為是thread private memory或者叫local memory。多數情況下每個線程會用多少local memory也是確定的。不過，如果出現一些把local memory當stack使用的復雜遞歸操作，可能造成local memory的大小在編譯期未知。這種情況編譯器會報warning，但是也能運行。不過local memory有最大尺寸限制，當前是每個線程最多512KB。
順序執行：亂序執行是CPU提高CPI的一個重要途徑，但亂序執行無論是設計復雜度還是運行控制的開銷都很大。CPU的亂序執行可以把一些不相關的任務提前（相關的也可以亂序，但要求順序提交），從而提高指令并行度，降低延遲。而GPU主要通過Warp切換的邏輯保持功能單元的吞吐處于高效利用狀態，這樣總體性能對單個warp內是否stall就不太敏感。
雖然GPU一般是順序執行，但指令之間不相互依賴的時候，可以連續發射而不用等待前一條指令完成。在理想的情況下，一個warp就可以把指令吞吐用滿。當然，實際程序還是會不可避免出現stall（比如branch），這時就需要靠TLP來隱藏這部分延遲。
顯式解決依賴：既然是順序執行，但同時又可以連續發射，那怎么保證不出現數據冒險呢？NV GPU現在主要有兩類方式：第一種是固定latency的指令，通過調節control codes中的stall count，或者插入其他無關指令，保證下一條相關指令發射前其輸入已經就位；第二種是不固定latency的指令，就需要通過顯式的設置和等待scoreboard來保證結果已經可用。在x86的CPU中，memory結果的可見性是通過緩存的一致性來控制的，這樣read-after-write之類的組合可以通過cache的可見性來保證，但多線程的情況也需要通過coherence和memory consistency model來保證。GPU本身運行就是多線程的，同一個warp內也是通過scoreboard來保證次序。但多個warp之間，GPU也需要維護相應的coherence和memory consistency model，具體大家可以參考PTX文檔: Memory Consistency Model。
當然這個邏輯雖然是這么設計的，估計偶爾也會有出bug的時候。Maxwell和之前的架構偶爾能看見編譯器往程序內插一些NOP。大概就是硬件上有問題，靠編譯器來強行修補。Turing上似乎已經比較少見了。
指令設計中的一些原則與思路
指令長度的問題
至少從Kepler開始，SASS就是定長的了。Kepler每條指令64bit，每8條指令含一條control code，之后的Maxwell、Pascal每條指令還是64bit，但是control code變成了每4條指令一條。Volta、Turing、Ampere都是每條指令128bit，每條都自帶control code。
定長和變長有什么差別呢？一般講體系結構的書上會用RISC與CISC來做對比，因為一般CISC指令集是變長的，比如x86。而RISC則通常是定長的。定長的好處之一是解碼器可以提前解碼，且一般解碼開銷小。因為首先指令等長，每個指令的范圍是確定的，指令定界不依賴于前一個指令的解碼。其次RISC在編碼的時候一般各個域的位置和長度比較整齊，解碼相對說來自然也更簡單一些。而變長則只能順序解碼，不得到前一條指令長度，后一條指令就不知道從哪里開始，當然就無法解碼。同時由于長度是變化的，每個操作數的類型和內容都可能會變化，解碼也就更繁瑣一些。不過變長有一個好處就是可以壓縮一些常用指令的長度，從而減少程序大小。而定長就沒有這種方式。
不過，現在很多架構和指令集設計都漸漸趨同。比如x86雖然變長（從一個Byte到十幾個Byte），但它解碼前會先定界，然后可以經過預解碼變成一系列的Micro Operation，這樣真正執行的時候也類似RISC。而RISC也不都是定長的，比如ARM的Thumb模式可以混用16bit和32bit的指令，RISC-V也可以加長指令實現特定的功能擴展。
對于GPU來講，我覺得長指令還是有很大的好處的。我的理由如下： - 指令Cache是GPU中命中率最高的Cache之一（constant cache也許是真的No.1）。一個warp 32線程已經讓指令解碼和調度之類的開銷大大均攤了。而同時大部分代碼都會被成千上萬個warp運行，這個開銷還可以被平攤得更小。當然，每個指令變長會加重指令Cache的容量和吞吐負擔，但它的格式也可以做得更整齊更有規律，從而大大簡化解碼過程。像x86這種編碼復雜的指令集，連解碼都會添加相應的cache（Micro-Op Cache）。而只要指令夠長夠整齊，解碼就可以做得夠簡單。這高ICache命中率的GPU來說，總體來講還是賺的。或者你可以認為，x86的Micro-op cache主要利好循環這種高重用代碼，而GPU每條代碼都與循環類似，那不如直接就用解碼后的指令做輸入。
- 前面提過指令執行的顆粒度。由于GPU流水線和發射邏輯的限制，每條指令都有基礎開銷（至少占用一次Dispatch Unit和相應Port），那在同一條指令中塞進更多信息，就成為減少指令基礎開銷的重要手段。指令越長，能編碼的信息就越多，比如指令內嵌的立即數就可以支持更長。典型的如ALU的立即數操作數，load/store的offset，branch的跳轉目標等。同時，多數指令也可以支持更復雜的modifier，從而表述能力和可編程性都更強。同時，更整齊規范的格式也更有利于編譯器做優化，性能的可及性更好。
- 指令夠長，就可以加入更多的調度和控制信息，從而簡化硬件自主判斷，拓寬指令觀察視野。Control code就是一個典型的例子。大部分的Control code其實是編譯期信息，沒有control code，其中很多依賴的判斷就需要專門的硬件邏輯去檢測和處理。而受限于硬件實現的復雜度，其視野肯定不如編譯時的全局視野寬。因此，通過這些控制信息，可以更好的利用編譯器在編譯時獲得的先驗知識，減少硬件自主分析和處理，從而簡化硬件設計，提高面積利用率的同時也簡化設計和驗證過程。
Volta開始的128bit指令，我覺得還是一個很有意義的嘗試。也許有人覺得這個也許可以做點內存壓縮，不過我感覺意義不是太大，壓縮和解壓其實也算是編碼解碼過程，尋址還會增加負擔，未必便宜。
指令集設計與ILP的一些相關性
GPU兩個最重要的并行邏輯，ILP（Instruction Level Parallelism）和TLP（Thread Level Parallelism，TLP有時在CPU語境下也代指Task Level Parallelism，兩者還是有所區別），兩者在隱藏延遲中都有重要作用。ILP的邏輯主要是靠前一條指令不需要執行完成就能發射下一條無關指令，而TLP則是通過warp之間切換來隱藏延遲。從另一個角度講，ILP和TLP都可以增加可發射指令的數目，盡量減少功能單元的閑置，從而提高硬件利用效率。
ILP是線程內（更準確的說是Warp內）的并行邏輯，影響ILP的主要因素有兩種，一是指令之間的依賴性，二是指令的資源競爭或沖突。依賴分顯式和隱式。顯式依賴主要是數據的相關性，隱式依賴則與資源競爭很相似，主要是兩個指令都要使用某個特定含義的公共資源。典型的顯式依賴如：
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兩條指令沒有數據依賴，仍然不能把2和3互換以隱藏1-2和3-4之間的延遲。而在Turing里這種指令已經可以顯式的用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。
這種相對更獨立的指令集設計其實有點類似函數式編程：操作專門carry寄存器可以看做是stateful的操作，改成可編程的Predicate后就成為只與輸入輸出有關的stateless操作，不改變機器狀態。也可以認為是所有需要用carry寄存器做輸入輸出的指令，都需要被序列化。通過與當前機器狀態解耦，獲得更大的指令調度自由度，編譯器的后端優化也會更加方便。
NV GPU在很早就開始有意識的淘汰這種含隱式輸入輸出的指令。比如早期使用隱式棧的call、ret、break等等（當然其實這和ILP關系已經不太大了）。例如SM61中要return時需要先顯式的用PRET設置某個隱式的調用棧，然后直接用RET返回。顯然在RET時這個棧必須是對應當初PRET設置的值（中間能不能再進出棧沒仔細研究），否則就會出錯。而Turing直接使用帶GPR地址的指令進行操作，就消除了這種隱式棧的操作過程，減少了指令之間復雜依賴對編譯器的干擾。
// 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來控制如何做浮點數的rounding，是否做subnormal的flush（x87 FPU control register控制普通的FPU指令，MXCSR控制SSE指令）。但在SASS中，FFMA、FMUL、DFMA等浮點運算指令，每個指令都可以自主控制是否打開flush（使用FTZ modifier），如何做rounding（RM, RP, RZ, RN）。這就意味著每個指令可以自主決定當前指令的運行方式，而不用改變機器狀態。不同模式混用時，就不需要保存和恢復control register了。
那再再擴展一點，x86其實也有控制FP exception的寄存器，NV GPU里是怎么操作的呢？我好像沒看見，感覺是被去掉了。這還是一個挺值得思考的問題~
當然，也不是說所有的隱式都應該變成顯式。比如每個Warp都有一個隱式的active mask，用來標記當前warp中divergence的情況。active mask與指令predicate的“AND”會共同決定當前指令是否起作用。那把mask交由指令顯式操作有意義嗎？我覺得沒有，因為這沒有太多額外的可編程價值。首先divergence造成的mask變化只能被分支或分支同步指令修改，其他指令需要控制效果直接用predicate就可以了，沒必要操作mask。其次，這個操作本身是非常固化的，增加相關操作指令并不會帶來新功能，反而會增加指令負擔（相當于每次可能有divergence的分支時，都要顯式的保存和設置mask，遍歷不同divergence的分支時就更麻煩了）。因此，在有predicate的情況下，active mask還是做成隱式的比較合理。當然，volta后的Independent Thread Scheduling也要自主控制和依賴內部的mask狀態，就更沒法做成顯式的了。不過，雖然mask是隱式的，但divergence后重新converge一般是顯式的（通過BSSY和BSYNC指令），否則程序就不知道應該在哪個點join了。
復合操作與附加操作
前面也提到了每個指令都有基礎開銷，那盡量在一個指令里塞進更多事情以分攤基礎開銷，就成為指令改進的一個重要方向。畢竟每個指令都要占用發射機會，但提高頻率很困難，也就是發射的總指令數有限，那就只好讓每個指令多做事。
比如一般的ALU指令可能有1~3個輸入操作數，為了盡量利用GPR的讀吞吐和增加每個指令的操作能力，SASS里其實有單指令多操作的情況。最常見的就是FMA（包括FFMA、DFMA、HFMA2和IMAD等），它可以一條指令同時計算乘和加兩個操作：d=a*b+c。Tensor Core的MMA類的指令也通常是3個輸入操作數。這里以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做操作數，不過這不是通常意義上的ALU指令，其實更接近memory指令一些，所以也沒列出來。然后是大部分非tensor類的整數和位操作指令都有Uniform datapath的對應版，這里也沒列。
這些多操作指令往往都需要更多的操作數。上一期我們講吞吐的時候也提到了，GPR的吞吐用不滿也就浪費了，所以一般ALU最多是3個輸入GPR。其實每cycle最多能讀4個，不過多少還是要留點余量給其他異步指令（如內存讀寫指令），否則搶占GPR的端口和搶占發射機會是一個效果。立即數和Predicate因為不占用GPR的吞吐，所以還可以額外加。比如LOP3.LUT R9, R15, R9, R16, 0xfe, !PT ;就是3個GPR+1立即數+1Predicate的輸入。于是，NV在設計指令集的時候又經常會在一些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里有一個函數叫__saturatef(float x)，它會把輸入clamp到[0,1]這個區間里。但是它并不占用一個指令，因為float指令自帶一個modifier叫.SAT，如FFMA.SAT R6, R2, R7, R0 ; 會自動對結果R6做saturate，并沒有額外開銷。類似的這種飽和操作在一些整數操作中也會出現（比如I2I.S16.S32.SAT R11, R11 ;表示把S32的輸入saturate到S16的范圍）。
復合操作和附加操作可以大大增強指令的表述能力，用更少的指令做更多的事情。但也要看到，它增加了指令的復雜度，對編譯器后端優化也提出了更多的要求。同時，在硬件設計上它也會有一些額外的開銷。如果這些操作不是特別常用，對指令運行開銷反而是負擔，所以多數還是需要一些功耗控制措施。比如IADD3其實多數時候都有一個是RZ，LOP3也通常只有兩個有效輸入。至于SETP后帶Predicate的情況，其實占比也不高。所以這些在具體設計和實施上，應該還是需要做相應的優化的。
立即數操作數和Constant Memory操作數
由于Volta前的SASS指令只有64bit，一般ALU的立即數操作數中只有19bit的編碼。這對于一些特定的指令是不太夠的，所以早期也有特定的32bit立即數的指令，如FADD32I，FMUL32I，MOV32I等等。不過這些指令最多只有2個輸入操作數，FFMA這種就沒法弄了。而計算很多函數值（比如sin）一般都是從高次多項式近似開始，需要計算大量給定系數的FMA，這些要用立即數就只能用mov先把32bit立即數寫入GPR，然后再FFMA。Volta后每個ALU包括FFMA都能用滿32bit立即數，就沒有這個問題了。
當然，實際Volta前的多項式逼近不是用立即數實現的，而是用Constant Memory。因為Constant Memory需要的編碼更少（幾bit的bank編碼+16bit的地址編碼），而且Constant memory還可以用在64bit的DFMA上，通用性更強。這也是constant memory對指令表達能力的一個重要貢獻。
Constant memory與立即數在做操作數時還有不少區別： 1. Constant memory運行有overhead，啟動kernel前需要初始化。而立即數是hardcode在編碼里，沒有額外overhead。 2. 立即數是編譯期常數，編譯時就必須知道值。Constant memory可以做運行期常數，也就是啟動kernel前才需要得到具體值，在運行前是可調的。 3. Constant memory有容量限制，當前應該是每個bank 64KB。立即數的總容量則是程序能寫多長就可以多大。 4. Constant memory運行期開銷會大一些，因為它畢竟是內存，需要相應的cache和load單元支持。 5. 立即數的編碼能力比較有限，當前128bit的指令也就32bit立即數。而Constant memory只要在指令格式中約定好，32bit和64bit都可以，將來擴展到128bit甚至更多也不是不行。
當前在CUDA程序里，編譯期常數是放在Constant memory還是做立即數是編譯器決定的。一般來講，32bit能精確表示的是立即數（比如float的所有數，double的1.0，2.5之類），需要64bit才能精確表示的會用constant memory。
注：傳運行期常數，直接用kernel的參數不也可以嗎？有什么區別呢？
其實很相似，但確實也有區別。因為kernel并沒有什么來自host的調用棧，kernel的參數其實也是用constant memory存儲的，在kernel啟動前會由驅動自動初始化。所以兩者從性質上很相似，但又有點差別。kernel參數的scope是當前kernel，每次啟動kernel理論上都要重新初始化參數區的constant memory，這其實也是kernel啟動的overhead之一（不太確定cuda Graph會不會做優化）。而用戶自己設置和copy的constant memory，scope至少是module（可以簡單認為是一個cu文件，或者更具體一點是cubin文件）。這就意味著同一個module的kernel都可見。所以如果同一個module內的kernel共用大量參數，且中間不會更改，那就只需要初始化一次（重復啟動同一個kernel也是同理）。
這甚至可以是一個極致優化的例子。把所有kernel的輸入參數都放在某個struct里，然后把struct復制到constant memory，就可以不用參數啟動這所有的kernel了。這應該會節省一點overhead（如果cuda graph不能做優化）。
PS：其實這個操作的意義主要在于接口的可擴展性。需要增加新參數的時候，只需要在struct里新加一個成員變量就行，不用去改函數參數和調用kernel的地方了。相信我，這個有時候真是很方便！你當然也可以在kernel參數里傳一個struct指針，這樣代碼量也相當。但是你需要自己把struct內容memcpy到device上，讀參數時還會增加內存負擔。因為constant在SM內有專門的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本身是算整數的d=a*b+c，類似浮點的FMA（前面也提到了，實際上Turing的IMAD和FFMA共用了功能單元）。單獨的IMAD就是三個32bit的計算，保留低32bit。IMAD.HI就是保留高32bit，IMAD.WIDE表示加數c是64bit（a、b還是32bit），輸出d也是64bit。這都算是常規操作。
后三個就有點意思了，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; 有什么區別呢？功能上兩者應該是一樣的，但功耗上可能不太一樣。因為IMAD會用到一個很長的乘法器，功耗會很大。但如果事先知道這個乘法器不用了，那硬件上就可以繞過或是用一些處理不觸發它，從而節省功耗。后面兩個指令也是類似，IMAD.SHL相當于b一定是2的冪次，IMAD.IADD相當于b一定是1，其實也是不需要完整的乘法器功能。不寫SHL或IADD不影響結果，但顯式的告訴功能單元，可以讓功能單元得到更明確具體的信息，從而進行一些優化。
那能不能讓硬件自己檢測呢？當然可以，但是未必劃算。首先，如果大多數IMAD指令都是不帶特殊性的操作數，那為了檢測這種case的功耗優化方案就會給硬件造成額外負擔。就像cache一樣，hit的時候自然是好，但加了cache往往也會劣化miss時的開銷。如果miss比率太高的話，加cache就是個負優化。同樣，在這里如果大部分情況執行路徑可以被優化，那自然可以省功耗。但如果多數情況其實優化不了，那就白白損失了優化檢測的這部分硬件開銷，變成了負優化。當然，增加這個modifier本身也是要看使用頻率的，如果極少用到，那其實也是降低了面積效率。這都需要對使用場景和規劃做準確的判斷。
這個邏輯其實也和control code有些相通之處，盡量把編譯期的信息融入到程序中，讓程序盡量明確運行邏輯。這其實也許體現了NV硬件設計中一個更上層的指導思想：軟件盡量具體地告訴硬件運行方式，而硬件則盡量減少自身的判斷和復雜監測邏輯，無腦的運行軟件的指示即可。這樣硬件設計盡量簡化，不但能節約面積開銷，測試、驗證復雜度都能降低。當然，這不是一件簡單的事情，想有效的安排和傳遞各種編譯信息，肯定需要軟硬件協同設計。這對軟硬件的架構規劃和協作，包括對應用、編譯器、指令集、硬件架構等統籌規劃，還是提出了很高的要求，很體現功力。
這里IMAD做這些優化應該是和編譯器有確定配合的，編譯器在合適的時候會優先選擇這些模式。其實這些功能也可以有其他一些實現方式，但是多則惑，也沒有必要。開銷大致等價的實現中選一個就可以了。
一些具體的指令
指令集內容實在太多了，有很多東西我沒有靠譜的輸入，瞎猜也意思不大。我就簡單列一些SASS中我覺得還比較有意思的點，供大家參考。這里基本都是Turing的指令。
FMA,MUL, ADD 系列
FFMA是衡量GPU算力的標桿之一。從指令吞吐角度講，FFMA一般都是最高的那一組。FFMA功能上相當于FMUL+FADD（FFMA精度更好，因為只做了一次舍入），那FMUL與FADD與FFMA共用單元就很合理。不過另一個有意思的地方是，FADD的reuse cache用的slot是1、3，是不是暗示它其實是FFMA的套殼，只是把FFMA第二個操作數固定成1？當然，更省的辦法應該是有直接的短路邏輯。
DFMA, DMUL, DADD系列具有不定長的延遲，需要靠scoreboard去控制依賴。我覺得一是因為FP64相關單元配置變化很大，有的還是多個SMSP之間share，可能存在競爭，這樣延遲不可控。二是它的吞吐可能很小（有的卡是1/16），指令延遲可能會超過stall count能有效表示的范圍，所以還是用scoreboard比較靠譜。
這些都支持32bit立即數做操作數。D系列雖然每個操作數是64bit，但如果后32bit都是0，那也可以放進32bit立即數里，用的時候做padding就可以了（指數位還是按FP64來，只是截了尾數后32bit）。HFMA2、HMUL2、HADD2則相當于把它當成兩個16bit的立即數。
IMAD, LEA, IADD3
IMAD在前面已經介紹過一些了。這里可以稍微再補充一些相關的。
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 ;
前三個我們已經介紹過，為了控制乘法器的開銷，有相應的modifier去明確輸入。但是可以發現，加數為RZ時沒看到專門的modifier。IADD3其實多數情況下也只有兩個操作數，第三個是RZ的概率很大。也沒有看到有modifier去優化這個，也許單純檢測RZ開銷不大，當然也可能是加法器的開銷沒那么大，不是很敏感。
另外就是IMAD在接受carry輸入時，有的前兩個輸入也都是RZ，但是這次就沒見到相應的modifier。不知道是沒設置還是沒有，可能是加了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數的任意按位邏輯運算。其實邏輯也很簡單，把這個映射關系看成一個函數d=F(a, b, c)，abcd都是bool值。輸入狀態有2^3=8種可能，那每個F函數都可以用這8個輸入時的輸出來完全表示。用8bit的編碼（立即數查找表，immeLUT）就可以指定這樣一個邏輯函數。更具體的用法和介紹大家可以參考PTX文檔。
LOP3 ImmeLUT
粗看起來好像挺厲害的，其實多數時候也沒有什么大用。我曾經搜過官方庫里的所有LOP3指令，看到的情況是：絕大部分情況下第三個操作數是RZ，也就是說3輸入的按位邏輯函數其實還是挺少見的。如果真的要找場景，可能密碼學或者挖礦之類會有些不錯的應用。
不過細品一下，其實還是有一些妙用的。比如我有兩個32bit數a，b，我需要把a的某些bit和b的bit拼在一起（對位替換），怎么弄呢？這其實就是一個簡單的按位邏輯函數c[i] ? a[i] : b[i]。這樣lop3就能用上。那什么時候會用到這種功能呢？其實還是有一些。比如copysign，需要把a的符號位替換掉b的符號位。再比如浮點數里有些指數、尾數之類的操作，也是有一些用處的。
更具體一點的例子，比如y= (x>=0) ? 1 : -1，這相當于就是把x的符號位復制給整數1。不過…… 這個一條LOP3指令做不到，因為這里c我們已經用了立即數（0x70000000，也就是選中第一個符號位bit），a或b就不能再用立即數了，等于說這里需要先把1或0x70000000移到GPR里，才能用LOP3來操作。
MUFU
MUFU是SASS中計算各種超越函數的指令。數學上，超越函數是相對代數函數（有限次加、減、乘、除、開方等組合）而言。但硬件上不能用多項式表述的好像都歸在超越函數里了。多項式求值一般就是乘和加，不需要其他的指令。而超越函數，就需要通過一些其他指令或是軟件逼近來實現。有的地方也叫特殊函數，我感覺是不太合理，就這些初等函數怎么也談不上很特殊吧……
常見的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 ;
其實數學函數求值這個事情還是有很多可聊的點。
首先，這些指令都是給出近似值。對精度要求高的話還是需要調用相應的數學函數庫再做軟件實現。一般說來，有的數學函數可以迭代更新（比如求倒數reciprocal，RCP），那就可以從近似值開始用不動點迭代得到更精確的值。另一種是可以做argument的range reduction，把全范圍的參數縮放到能比較精確計算的范圍內。比如對于x具有二進制形式x=a2^e, log2(x) = log2(a2^e)=e + log2(a)。那只需要精確計算log2(a)（其中a在1~2之間）就可以得到精度很高的解，而MUFU.LG2在某些特定范圍內是可以滿足精度要求的。還有一種就是sin和cos這種，它沒有迭代形式，也沒有特別好的手段做規約，那就先規約到一個不太大的范圍（比如正負pi/4），然后用一個高次多項式去近似（不是泰勒展開，大家有興趣可以去研究一下Remez的minimax方法）。所以MUFU.SIN只是在近似計算中出現，精確計算中可能用不著它。
當然，軟件實現上還是有很多自由度，這里只是給出了一種選擇方式。NV的數學函數庫不直接開源，但是有相應的LLVM的bitcode（NVVM目錄下的libdevice.*.bc），可讀性還比較強，可以參考。
從這些實現可以看到，大量輔助計算仍然是用加、乘這種初級運算完成的，所以即使是使用了大量數學函數，主要的計算量也在FMA、MUL和ADD這種簡單指令。MUFU還是占比比較小。如果使用了內置函數，則相對占比高一些，但多數時候還是需要相應的簡單指令的配合。而且有些常用函數其實并沒有加速實現（比如sqrt，即使是intrinsic的__fsqrt_rn這種實現也非常復雜）。這里面肯定是有所考慮的。
MUFU的64bit版本（RCP64H和RSQ64H）其實只用了一個GPR輸入，估計覺得反正是近似，多后面32bit意義也不大。
FSETP
FSETP本身也沒有什么特殊的，就是根據一些浮點的比較操作設置一個bool變量。比如：
FSETP.GT.AND P1, PT, R1, R2, PT ;
FSETP.GEU.AND P1, PT, R1, RZ, PT ;
第一個是比較R1是否大于（greater than，GT）R2，然后順帶把結果與另一predicate做了與（AND）。GE是greater or equal，那GEU是什么呢？這其實是浮點比較中的一個特例。PTX把這個叫unordered floating-point comparisons，專門為了NAN的特殊形式。
如果熟悉IEEE 754，就知道NAN如果出現在比較浮點運算（a op b）中，那結果永遠是false，甚至x為nan時x==x都是false（甚至能用這個來判斷x是不是nan……）。所以PTX或SASS中有兩類浮點比較，一類是ordered，有nan就為false，一類是unordered，變成只要有nan就為true。如果兩個輸入都不是nan，那兩者就一樣。
這個邏輯有什么用呢？一些編譯器在處理比較運算的時候，會選擇一個canonical的形式（比如把a >= b統一改成 b<a，這樣就不用處理>=這種操作了），而對浮點操作這個變化是不正確的，就是因為有nan的存在。但是ordered a>=b 與 unordered b<a 兩者是等價的。
浮點的比較操作里還是有不少坑，和整數的比較還是差別很大。除了nan的不等于自己，還有一個就是+0和-0的二進制表示不一樣，但卻是相等的。denormal的情況我沒有仔細研究，要是會被flush的話，這里面也有很多需要注意的問題。
LDG/STG, LDS/STS, LDL/STL
內存相關指令一般都比較復雜，特別是global操作，modifier特別多。這部分其實已經超出指令集設計本身的范疇，更多的是consistency model的問題。這只討論指令輸入輸出中的一些點。
比如說LDG的有兩個使用UR做base的形式：
LDG.E.SYS R16, [R10.64+UR10+0x200] ;
LDG.E.CONSTANT.SYS R54, [R32.U32+UR4] ;
第一條中R10.64應該表示它是個64bit的值，如果是32bit會用第二行的R32.U32這種形式。而后面的UR4，一般都是64bit。用32bit的好處就是一些簡單的地址運算可以用32bit的運算得到，比每次都做64bit操作要更劃算。而如果是warp內uniform的base變化，可以直接改UR，UDP的計算與一般的ALU也是獨立的。這也算是一種減少計算強度的方式。
LDG/STG也有一些控制cache的邏輯，比如帶CONSTANT的modifier表示會在read-only L1中緩存。當然，這個具體操作其實每代架構都會有所不同，但能用constant的話沒理由不用。
Shared Memory里有一個有意思的內置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，基址也常常是編譯期常數（靜態分配）。所以地址計算往往可以表示為immeBase + index * size。而元素大小常常是2的冪次，這樣就可以省去預先計算乘法或是移位，直接讓地址單元進行處理。這也算是減少計算強度的方法。
不過其實LDG/STG/LDL/STL應該也可以用這個邏輯，特別是local memory，不知道為什么沒有。
BAR
BAR指令就是barrier，更確切的說是synchronize barrier，就是通常block內用來同步的__syncthreads()函數。CUDA中叫barrier的術語挺多的，比如dependency barrier主要指DEPBAR指令，是用來等scoreboard的。memory barrier應該是相當于memory fence，但是現在好像有專門的API操作。還有一種convergence barrier就比較底層了，是用來處理warp divergence的。
CUDA C一般只用一個barrier（就是BAR.SYNC 0x0后面的那個0x0），但實際上PTX或SASS里每個block最多可以用16個barrier，每次同步可以選擇同步到某個barrier。這樣同一個block的線程可以分塊同步到不同的barrier上去（BAR有warp計數的參數）。另外，BAR還支持arrive、sync模式，可以搭配出類似生產者-消費者的模型，具體可以參考PTX文檔。這個功能沒有C API，但是可以用inline PTX其實風險也不大。這其實也算是一個降低瓶頸資源壓力提高并行度的方法，與前面提到的用predicate做carry，還有跳轉用顯式GPR代替隱式棧，應該算是一類思路。
CUDA 11之后有一個新功能叫Asynchronous Barrier。從實現上更偏軟件，但功能上沒看出它比BAR.arrive與BAR.sync的組合有什么更特別的地方。也許是scope更靈活？或是更適合CPU軟件移植和編譯器實現？
BAR不僅可以實現同步，還可以順帶計數或是做reduction，這個是可以從CUDA C調用的（參考）。從功能上講，用戶自己用shared memory實現應該也可以，不過既然是順帶的，多少還是省點用戶的事。具體硬件實現上便不便宜，就不得而知了。
結語
指令集的設計其實是一個挺玄學的問題。好的指令集設計與好的產品之間，其實還隔著非常多其他因素。功能貼合市場需求，性能滿足要求，其實就具備了成為好產品的條件。而指令集設計本身依附于產品提供的功能和性能，它更多的是一個接口的角色：連接用戶與設備，讓功能和性能更容易發揮。但我覺得指令集并不是決定性因素。x86指令集現在常被吐糟是過時設計，即使它已經不太強勢，但生命力仍然頑強。我覺得也不是光靠“兼容性”和“wintel壟斷”可以解釋得通的，至少它當前具有的功能和性能，一定是多年市場選擇和無數技術迭代的結果。當然，好的指令集設計，可以讓一個架構更通用，性能可及性更好，硬件效率更高，從而更有生命力和競爭力。
二．CUDA微架構與指令集-指令發射與warp調度
CUDA指令的發射和warp調度問題。
? 指令發射的基本邏輯，主要是指令發射需要滿足的條件，幾代架構發射指令的一些簡單描述等等。
? control codes，主要是它與指令發射和warp調度的關系。
? warp Scheduler的功能。
? 峰值算力的計算方法和達成條件。
指令發射的基本邏輯
首先，簡單回顧一下CUDA程序的等級結構。每個Kernel有一個grid，下面有若干個block，每個block有若干個warp。同一個block的warp只能在同一個SM上運行，但是同一SM可以可以容納來自不同block甚至不同grid的若干個warp。我們這里要聊的指令發射和warp調度的問題，就是指同一個SM內同一個warp或是不同warp的指令之間是按照什么邏輯來調度運行的。
指令發射的一些基本邏輯：
- 每個指令都需要有對應的功能單元（Functional Unit）來執行。比如執行整數指令的單元，執行浮點運算指令的浮點單元，執行跳轉的分支單元等等。功能單元的個數決定了這種指令的極限發射帶寬（在沒有其他資源沖突時）。
- 每個指令都要dispatch unit經由dispatch port進行發射。不同的功能單元可能會共用dispatch port，這就意味著這些功能單元的指令需要通過競爭來獲得發射機會。不同的架構dispatch port的數目和與功能單元分配情況會有一些差別。
- 有些指令由于功能單元少，需要經由同一個dispatch port發射多次，這樣dispatch port是一直占著的，期間也不能發射其他指令。比較典型的是F64指令和MUFU特殊函數指令。
- 每個指令能否發射還要滿足相應的依賴關系和資源需求。比如指令LDG.E R6, [R2] ;首先需要等待之前寫入R[2:3]的指令完成，其次需要當前memory IO的queue還有空位，否則指令也無法下發。還有一些指令可能有conflict的情況，比如shared memory的bank conflict，register的bank conflict，atomic的地址conflict，constant memory在同一warp內地址不統一的conflict等等，這些都有可能導致指令re-issue（甚至cache miss也可能導致指令replay）。這些情況會不會重復占用dispatch port發射帶寬我暫時還沒仔細研究。
- 在有多個warp滿足發射條件的情況下，由于資源有限，需要排隊等待發射，warp scheduler會根據一定的策略來選擇其中的一個warp進行指令發射。
- 當前CUDA的所有架構都沒有亂序執行（Out of order），意味著每個warp的指令一定是按照運行順序來發射的。當然，有的架構支持dual-issue，這樣可以有兩個連續的指令同時發射，前提是兩者不相互依賴，而且有相應的空余資源（比如功能單元）供指令運行（對kepler來說不一定是不同類的功能單元，后面會具體分析）。另外一個顯而易見的要求是雙發射的第一個指令不能是分支或跳轉指令。
接著我們要簡單講講Kepler，Maxwell，Turing三代架構的指令發射和warp調度邏輯。先看看這三個微架構典型chip的SM示意圖（圖來自各自whitepaper，注意SM的名稱有點變化）：
簡單對比一下其中的warp scheduler和dispatch unit：
每個warp scheduler每cycle可以選中一個warp，每個dispatch unit每cycle可以issue一個指令。幾個架構的區別：
- Kepler的SMX有192個core，但是core和warp沒有固定的從屬關系，相當于每個warp都可以運行在不同的32core組上（？這個需要進一步確認）。192=32*6，所以每個cycle最少需要發6個指令才能填滿所有core。可是4個warp scheduler只能選中4個warp。所以Kepler必須依賴dual issue才能填滿所有cuda core，而且由于core多且與warp沒有對應關系，kepler的dual issue不一定是發給兩個不同的功能單元，兩個整數、兩個F32或是混合之類的搭配應該也是可以的，關鍵是要有足夠多的空閑core。每個warp scheduler配了兩個dispatch unit，共8個，而發射帶寬填滿cuda core只要6個就夠了，多出來的2個可以用來雙發射一些load/store或是branch之類的指令。
- Maxwell開始，SM的資源就做了明確的分區，每個warp都會從屬于某個分區，各分區之間有些功能單元（比如cuda core和F64單元，SFU）是不共享的。Maxwell的SMM有128個core，分成4組。每組有一個warp scheduler，2個dispatch unit，配32個CUDA core。這樣每cycle發一個ALU指令就可以填滿cuda core了，多出來的dispatch unit可以發射其他memory或是branch指令。由于功能單元做了分區，沒有冗余了，這樣Maxwell的dual-issue就不能發給同樣的功能單元了。
- Turing的SM的core數減半，變成64個，但還是分成4個區，每個區16個core，配一個warp scheduler和一個dispatch unit。這樣兩個cycle發一個指令就足夠填滿所有core了，另一個cycle就可以用來發射別的指令。所以Turing就沒有dual-issue的必要了。從Volta開始，NV把整數和一些浮點數的pipe分開，使用不同的dispatch port。這樣，一些整數指令和浮點數指令就可以不用競爭發射機會了。一般整數指令用途很廣，即使是浮點運算為主的程序，也仍然需要整數指令進行一些地址和輔助運算。因此，把這兩者分開對性能還是很有一些幫助的。但是，這里還是要澄清一下，不是所有的浮點和整數都是分開的。這里貼一張Hotchips里NV介紹Turing的圖（圖上沒畫F64，應該是和Tensor Core、F16在一塊）：
Volta和Turing的F64、F16和Tensor Core都是用的同一個dispatch port。然后F32一組（IMAD也放在這組，大概是要共享mantissa的那個乘法器），MUFU（就是特殊函數指令），其他ALU（包括整數算術運算和移位、位運算等等，但是不包括IMAD）一組，然后memory指令一組，branch指令一組，然后Turing的uniform datapath的指令一組，這些都各自有dispatch port。不同組的可以隔一個cycle發射，同組的就要看功能單元數目，最少是隔2個cycle。
Control codes
從Kepler開始，指令中就帶有Control codes。Kepler的架構略顯久遠，格式又不太一樣，信息也不太全。Maxwell后control codes的功能更加豐富了，而且在Scott Gray的文章中非常具體詳盡的描述了各個域的含義和功能。Pascal的指令集與Maxwell基本一致，control codes也沒什么變化。Volta和Turing，包括最新的Ampere，只是把control codes編碼到每條指令中，具體內容和bit對應關系其實也沒有變，所以我這也不做架構的區分。這里主要復述一下Scott Gray的表述，但主要側重它與指令發射和warp調度的關系。
這里我以Turing為例，選了一段程序，把對應的control codes 寫在前面（這里我用的格式與Scott有點區別，主要是看起來更方便）：
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，對應4個slot，有reuse就寫R，沒有就”-"），Wait Dependency Barrier（6bit，B+6個數，有等待就寫上對應的barrier號0-5，否則寫“-”），Read Dependency Barrier（3bit，R+設置的barrier號，不設置寫“-”），Write Dependency Barrier（3bit，W+設置的barrier號），Yield Hint Flag（1bit，“Y”表示yield，否則寫“-”），Stall Count（4bit，S+十進制的stall的cycle數）。
這里與scott的標記的區別一是顯式的寫了reuse，而不是只寫在后面匯編文本里。這樣看起來更直觀且容易看到slot的對應關系，因為有的predicate夾在中間容易搞混。二是wait barrier的每個bit我拆開了，這樣肉眼更容易與前面設置barrier的R#或是W#對應上。然后就是每個域前加了提示字符，BRWS之類，畢竟有的時候看久了也容易恍惚，這樣區分比較明顯，不容易搞混。
【更正1】：這里包括后文說的Dependency Barrier也許更合理的稱呼是Scoreboard，在profiler里有short scoreboard和long scoreboard之分，應該指的就是這個。
【更正2】：scoreboard的編號我調整過一次。由于最開始我沿用了scott的1-6的編號方式，但后來我還是調整為更符合原語境的0-5。首先這是原編碼的值，其次DEPBAR后會顯式的接SB0~SB5這種數，所以后來我就統一改為0-5。
下面具體介紹control codes每個域的含義。
Register Reuse Cache
Register Reuse Cache有4bit。每個指令有4個slot，每個register的source operand的位置對應一個slot（predicate好像不算）。我暫時還沒碰到4個source operand的指令，所以有一個bit好像一直沒用到。Reuse的用法：如果當前指令某個slot的register還會被下一個指令的同一個slot讀取，那就可以reuse當前指令讀取到的register內容。Reuse的作用主要就是減少GPR的讀取，一來可以減少register bank conflict，二來應該也能省一點功耗。比如前面代碼中Line6的IADD3的第一個源操作數是R113，而Line7的第一個源操作數也是R113，所以可以reuse。同理，line7、line8同樣位置的R113都可以reuse。但是line8的R0為什么reuse呢？我也沒搞懂。Reuse是唯一在官方反匯編中出現的control codes，但也有很多疑點。
我的幾個猜測：首先reuse是個cache，某種意義上是個hint，也就是說就算set了，reuse不了應該也不至于出錯。第二，reuse cache的位置應該是位于所謂的operand collector。這個應該是很多功能單元公用的，所以不太可能是同種功能單元的指令才能reuse。另外，load store類的指令沒看到reuse，好像是不用collector。第三，如果切換到別的warp，register是不同的空間，那reuse cache就失效了。所以reuse應該是需要當前warp的指令連續發射。第四，reuse既然是register的cache，那原來register的位置出現了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已經被Line2的R4用過了，應該是沒法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，編譯器做了相應的指令調度，但是reuse卻沒有重新生成。所以我懷疑這個是編譯器的bug，存疑中。
Wait Dependency Barrier
Wait Dependency Barrier有6bit，每個bit表示是否需要等待對應的dependency barrier。每個線程有6個dependency barrier，每個barrier都可以被后面的Read或Write操作設置上。設置wait dependency barrier是等待依賴的其中一種方式。SASS里面還有一個對應的指令，如DEPBAR.LE SB0, 0x0, {2,1} ;，兩者自由度不太一樣。control codes里設置的barrier只能是bool形式，要么dependency resolved，要么就是not ready。而DEPBAR可以等待有計數的barrier。比如發了8個memory指令，其實設置的是同一個dependency barrier，每發出一個計數加1，每回來一個計數減1。這樣DEPBAR可以等計數降到6就說明前面2個指令已經到位了，對應的load結果就能用了。而如果用control codes來等待，就只能等計數降到0，也就是8個都ready才行。
dependency barrier有一個和分支指令強相關的地方，比如不確定的跳轉指令（帶predicate的BRA，或是BRX這種指令）需要等待當前所有已設置的dependency barrier都到齊才行。否則后面多個分支的代碼可能用到這個barrier，但又不一定都會等待。這個可能是編譯器在處理上有一些圖方便的地方。有些情況是可以把wait后移到對應使用指令上的，這樣延遲更容易被隱藏。只是有時候編譯器拿不到足夠的信息，為保證正確性就統一在跳轉的時候等了。
Read Dependency Barrier
Read dependency barrier有3bit，表示需要設置的6個barrier中對應的索引（0~5，對應barrier 1-6，如果不需要設置barrier，就設置為0b111）。Read dependency barrier主要是一些指令不會在一開始就把所有操作數讀進去，所以需要hold住GPR的值，防止后面的指令在它讀取其內容之前把GPR改掉。使用Read dependency barrier的主要是memory類的指令，但是一些轉換指令如F2I/I2F之類好像偶爾也能見到。
Write Dependency Barrier
Write dependency barrier與read dependency很類似，也是3bit，后面跟barrier索引。注意Read和Write兩者用的dependency barrier資源是一樣的，也都是上面wait的那6個。Write dependency barrier比較好理解，就是某個指令要把操作結果保存到某個GPR或是predicate中，使用barrier進行保序可以防止出現data race。不過這主要針對的是不定長latency指令。如果一個指令的latency是確定的（或者有不太長的上限），那用后面提到的stall cycle停足夠長時間就可以保證沒有race。
Yield Hint Flag
Yield hint flag是1bit。如果Yield，就表示下一個cycle會優先發射其他warp的指令。前面聊reuse的時候我們也說了，reuse是需要連續發射同一個warp的指令的。所以reuse和yield是不會聯用的。另外，SASS有一個專門的指令YIELD，感覺上是一樣的功能（【注】：經評論區提醒，功能還是不一樣的，具體請移步評論區）。Yield的存在仿佛暗示warp scheduler不是round robin的選擇warp，而是傾向于一直往同樣的warp里發射指令（如果可以一直發射的話，有stall肯定就盡量切走了）。但是這個東西我也不太確定，沒仔細測過。如果真是這樣，yield的作用就是保持各個warp之間的進度均衡，否則在barrier之類的指令上會有較大的等待開銷。而且如果退出時間差很大，也會導致一些資源不能盡快回收以容納新的block。
【補充】：根據評論區所說，也許yield這個bit就是stall count的高位。只是假如這個bit不為0，那stall的cycle會>16，相當于warp被切換的概率也會大大增加。兩者之間的具體含義應該還和具體指令有關，不同類型的指令也許是不一樣的。這個問題我還沒有具體研究過。先存疑。
Stall Count
Stall count有4bit，表示當前指令后需要stall指令發射的cycle數，然后再決定是不是要繼續發射。這個cycle數受到極限發射帶寬的約束，很多時候可以用來反推功能單元的分組和數目。比如Maxwell下有雙發射，所以可以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的雙發射不限定功能單元，Turing架構沒有雙發射，就都沒有這種形式了。
我們再找個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] ;
前面我們已經講過，Turing的兩個周期發一個ALU指令就可以用滿cuda core了，所以同組ALU指令至少stall 2 cycle。如果兩個指令間只stall 1 cycle，說明這兩個指令應該分屬不同的功能單元（或者說不共用dispatch port），可以分開發射。如果stall的時間更長，說明其發射帶寬比較低，比如LDG/LDS/STG/STS這種memory指令都是4 cycle才能發一個。但是也有一些奇怪的例子：
[----: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居然是挺機關槍，每cycle都可以發。可是按programming guide中的Instruction Throughput表，I2F如果不涉及F64，是1/4的throughput。難道是有特別的queue做buffer？這個也沒太明白，有待進一步研究。
其他
Control codes一個比較容易忽視的問題是它與predicate是獨立的。也就是說不管加不加predicate，control codes的作用是不會改變的。因為本身control codes很多東西是編譯期決定的，如果按運行期的predicate來定是否啟用control codes，有些代碼的正確性就容易出問題。
Control codes是個很復雜的問題，有些東西真是純逆向了。我也沒仔細研究過。大家有興趣可以看看scott gray的maxas的實現，或許會有一些啟發。
Warp Scheduler
Warp scheduler的作用就是管理一系列的warp，在那些滿足條件的warp中選中一個來發射指令。就緒可以發射指令的warp就是eligible，不滿足發射條件的warp就是stalled。導致warp不能發射指令的原因有很多種。我根據NSight Visual Studio Edition中Issue Stall Reasons中的描述，大致搬運翻譯一下：
? Pipeline Busy：指令運行所需的功能單元正忙。
? Texture: Texture單元正忙，或者說已經下發的request過多。
? Constant：Constant cache的miss。一般說來，多數情況下constant cache的hit rate還是很高的，所以一般只會在第一次access的時候miss。
? Instruction Fetch：Instruction cache的miss。與constant cache的miss類似，一般只有第一次運行到的地方才容易miss。比如BRA新跳轉到的地方，或是Instruction cache的cache line的邊界處。
? Memory Throttle：有大量memory操作尚未完成，導致memory指令無法下發。可以通過合并memory transactions來緩解.
? Memory Dependency: 由于請求資源不可用或是滿載導致load/store無法執行，可以通過內存訪問對齊和改變access pattern來緩解。這個與memory throttle的細微差別我還沒仔細研究過。
? Synchronization：warp在等待同步指令，如cuda C里的_syncthreads()，對應SASS里的BAR指令。
? Execution Dependency：輸入依賴關系沒解決。簡單說，就是輸入值未就緒，就是在等control codes里的dependency barrier。單個warp內通過增加ILP可以減少依賴型stall。如果ILP不夠用，這個stall就會形成額外的latency，只能用TLP來隱藏了。
在Profiler里提供的Turing的performance counter里，還有兩個當前warp不能發射指令的原因:
? stall_not_selected: warp當前雖然eligible，eligible的warp超過一個，當前的未被選中，所以不能發射。
? stall_sleeping: 這個一般是用戶自己調用sleep功能讓該warp處于睡眠狀態。
貼一個Nsight里的stall reason統計圖：
Warp Scheduler的另一個關鍵功能是在eligible里選一個來發射。很多早期的書上都說選warp是round-robin，就是所謂的輪詢，發一個換一個。當前的幾代架構，至少maxwell之后，我覺得應該不是這個策略了。前面聊reuse，yield和stall的時候也提到了，如果是round-robin，這些東西都會顯得很奇怪了。所以我感覺它應該是比較aggresive的往同一個warp發射指令，除非stall了。當然，中間如果沒有yield，那stall 2個cycle的時候中間那個cycle能去發射別的warp嗎？這個我也有點迷，有機會再仔細研究下。
【補充】：我仔細再想了想，stall 2cycle的時候中間那個cycle應該是可以發射其他不用operand collector的單元，比如memory和branch。網上看到一些說法，有些S2R的指令好像也是走的memory的pipe，那應該也可以發射。這個應該是可以寫一個micro-benchmarking的case驗證一下，只是我手上沒有Turing的卡，也沒法測了。但是我感覺只要它不影響reuse的cache，那就應該可以抽空發射。這也是一個比較符合性能需求的模式。
Eligible的warp數是影響峰值性能的關鍵表征之一，如果每個cycle都至少有一個eligible warp，那功能單元基本就會處于滿載狀態，性能一般也會比較好。這也是occupancy真正起作用的方式。
關于峰值算力的問題
最后要略微展開聊一下峰值算力的問題。NV的GPU經常用CUDA Core這個詞來表示算力強弱，這其實是Int32/Float32功能單元的marketing叫法。算力評估常用的單位是FLOPS，表示FLoat OPerations per Second（Flops有時候也用作Flop的復數，注意鑒別）。對于F32而言，FADD/FMUL都是一個指令一個flop，FFMA一個指令同時算乘加所以是兩個flop。所以一般NV的GPU的F32峰值算力計算方法為：
SM數 * 每SM的Core數 * 2 * 運行頻率
最后結果常用GFlops或是TFlops表示。其中乘以“2”是因為FFMA是兩個flop。比如說Maxwell架構每個SM有128個CUDA core，每個SM每cycle可以發射128條Int32或Float32指令（兩種指令不能同時發射，所以是或）。Maxwell架構的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每秒可以發射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浮點峰值做例子，如果你的任務不需要浮點運算或是精度不是F32，這個值就意義不大，需要轉換成你需要的那個操作。現在AI處理器常常宣傳峰值是多少FLOPs，或是多少IOPs，一般也會限定是F32，F16或是I8之類。因為每種操作對應的指令是不一樣的，峰值當然也可能不一樣。頂級HPC計算卡F64一般是F32的一半，但消費級顯卡F64多數會有閹割。如果沒有TensorCore而用packed F16（把兩個F16塞到一個32bit GPR里同時運算），F16峰值性能通常是F32的兩倍。有TensorCore時則要另行計算，要看具體TensorCore的數目和指令帶寬，還有能不能和其他指令同時發射等。其實不看Tensor core的話，滿血版一般有：F64:F32:F16=1:2:4，正好與占用的GPR成反比，這個其實是與GPR的帶寬有很大的關聯的，一般滿血版的卡的功能單元配比就會盡量按極限的GPR帶寬來設計。
這里貼一個Ampere的white paper中V100和A100的幾種峰值性能對比：
要達到F32的峰值性能，需要滿載發射FFMA指令，這是很苛刻的條件。首先，其他與FFMA共用dispatch port的指令，每發射一個都會擠占FFMA的發射機會。其次，由于多數情況下數據要從memory中來，而memory操作比ALU慢很多，常常導致指令操作數無法就緒，從而有些周期沒有FFMA指令可發。同時還有其他一些overhead或是occupancy問題導致有些SM無法滿載，從而無法達到峰值。一般說來，實際應用中，較大尺寸的矩陣乘法（GEMM）是難得的能接近峰值性能的程序，有些實現能到98%峰值的效率。但多數實際應用效率都遠不及此，很多memory bound程序能到10%就很不錯了。超算TOP500排名中，多數HPL效率都是5060%左右，更接近實際應用的HPCG效率一般都在23%左右。雖有規模大導致的互聯開銷原因，但總體來講實際應用的峰值性能離極限值還是差距很大的。
每種類型的指令都有一個峰值性能，那是不是能同時達到呢？基本是不能。對于共用dispatch port就不說了，要相互競爭發射機會，發射一條這種指令就少發射一條那種指令，所以顯然不能同時到達峰值。如果是不同的dispatch port呢？理論上可以，但是實際上也會比較難。比如說Turing的F32和I32，首先I32的2IOP指令IMAD是和F32一伙的，相互競爭dispatch port，所以兩個不能同時到達峰值。剩下的IADD3或是LEA之類的指令理論上可以與F32的并行，倒是有機會沖一沖。只不過多數實際應用中很難做到這么好的運算配比，而且register的bank conflict之類應該也會大大限制這兩種指令的同時運行。另外，即使兩者配合完美，它還是需要省出一些發射帶寬給其他配套指令（比如memory load），不可能完全占滿。
本文將兩篇文章整理了一下。
參考文章鏈接如下：
https://zhuanlan.zhihu.com/p/391238629
https://zhuanlan.zhihu.com/p/166180054
總結
以上是生活随笔為你收集整理的GPU指令集技术分析的全部內容，希望文章能夠幫你解決所遇到的問題。
