VoltaTuring GPU Architecture - I&II

本文源自各種資料的拼湊和腦補，不保證所有主張的正確性，請自行斟酌。

I。 指令

I。I 指令構成

Pascal和Maxwell架構採用64位指令編碼，將每3條指令打包為Bundle，並附加一個64位的控制資訊，其中每21位用來編碼一條指令的控制資訊。

Volta和Turing架構拋棄了Instruction Bundle，使用128位的指令字長。這128位中，用91位 編碼

指令

，23位來編碼

控制資訊

，14位留用

［1］

。

V&T使用變長操作碼，長度範圍為10-13bits。運算元可以是

通用暫存器(RX)

，

uniform暫存器(URX)

，

特殊暫存器(SRX)

，

predicate暫存器(PX)

，

常量記憶體(c[x][y])

或

立即數

。

在23位的

控制資訊

中，總共分為幾個部分：

2位留空。

4位的

複用標識

（Reuse flags）。

6位的

等待掩碼

（Wait barrier mask）。

3位的

讀屏障序號

（Read barrier index）和3位的

寫屏障序號

（Write barrier index）。

1位的

讓步標識

（Yield flag）。

4位的

阻滯位

（Stall cycles）。

複用標識：

在NV的指令集設計中，單指令最多擁有4個運算元。而從Maxwell開始，為了減少暫存器bank confict，為每個SP增設了4個register reuse cache

［2］

。透過將複用標識置為1，可以指示對應的運算元暫存器值將會被送往對應的暫存器複用快取。下一條指令如果需要使用該暫存器的值，則無需再從暫存器中取數，而是直接從reuse cache中取數。

如圖中所示，指令0x0050標識了 複用R6、R11暫存器的值，因此SP在執行該指令時，將把這兩個暫存器的值送reuse cache，而指令0x0060的對應位置運算元剛好是R6和R11，這時就可以直接從cache中取數而不必讀register。

這一特性為解決register bank conflict提供了一種新途徑。

同時也有助於減少資料競爭。

等待掩碼、讀/寫屏障序號：

用掩碼的形式指出該條指令必須等待某些訊號。雖然大部分指令能在確定的週期內完成，故而能夠被彙編器靜態排程，但訪存指令往往會帶來可變的延遲，造成無法避免的

動態指令排程

。Volta/Turing使用依賴屏障機制來規避資料冒險。

RAW：當一條可變延遲指令（如0x0070）需要寫某一暫存器（如R2），則彙編器為該指令分配一個可用的屏障號碼，並將該號碼寫入讀屏障序號中。後續讀暫存器指令（如0x0090）若需要訪問暫存器R2，則彙編器將該指令的等待掩碼的對應位置為1。若0x0070指令未完成，則對應讀屏障仍存在，與0x0090指令的等待掩碼做與操作後不為0，則0x0090指令將被阻塞，實現執行時指令排程，防止RAW競爭發生。

WAR：當一條可變延遲指令（如0x00e0）需要將暫存器（R9）內容寫入記憶體時，必須保證該暫存器的值在指令完成前不被更改。則彙編器為該指令分配一個可用的屏障號碼，並將該號碼寫入寫屏障序號中。後續寫暫存器指令若需要寫暫存器R9，則彙編器將該指令的等待掩碼的對應位置為1。若0x00e0指令未完成，則對應寫屏障仍存在，該寫暫存器指令將被阻塞，防止WAR競爭發生。

由於等待掩碼以mask的形式產生作用，故而一條指令可以等待多個屏障（如當運算元的兩個源暫存器都有RAW）。

讓步標識：

該標識位標識了期望的執行緒排程行為。當該位為1時，指示排程器，期望在發射該指令後，繼續發射當前warp下條指令而不切換warp。當該位為0時，指示排程器期望在發射該條指令後切換warp。

但排程器在執行時的實際行為不完全受該標識位約束。

阻滯位：

用4個bits來編碼阻滯週期數。當該部分為0b0000時，排程器會直接發射下一條指令。若該部分不全為0，則指示排程器等待相應的週期數（最多15個週期）後，繼續發射下一條指令。

I。II 指令集

V&T指令集有以下幾種型別的指令

［3］

：

浮點指令。又可分為16位，32位和64位指令。

整數指令。包括算數與邏輯運算和位操作。

型別轉換指令。

資料傳輸指令。包括暫存器間傳輸以及shuffle等。

斷言指令。

讀寫指令。包括記憶體讀寫，原子操作，快取控制，記憶體屏障等。

Uniform Datapath Instructions。不知道是什麼。

紋理指令。讀取紋理，應用filter等。

表面指令。讀寫表面資料等。

控制指令。呼叫，返回，斷點，跳轉，終止等等。

其他指令。這個很雜，我也看不懂幾個。看得懂的比如有：讀特殊暫存器，NOP，設定local memory的基址，設定協作執行緒組ID（wtf，這玩意居然允許使用者更改嗎）。

II。 執行緒排程模型 - 通用計算

GPU是一種能夠並行執行大量執行緒的計算裝置。大量的執行緒同時執行相同的程式，並處理不同的資料。為了更加高效的執行儘可能多的執行緒，NV提出了

CTA(Cooperative thread arrays,執行緒協作組)

的概念，僅允許同一CTA的執行緒之間相互通訊，分屬不同CTA的執行緒將無法進行直接通訊。

NV的GPU架構由可增減的一組

SM(Streaming Multiprocessors,流式多處理器)

構成，每個SM之間不會發生直接互動，故而可以任意增減SM的數量而不對架構做任何改變，也可以支援多塊GPU協同工作。當CPU程式發起一項任務後，

執行緒將會按CTA被分配給不同的SM執行

。

當一個SM完成當前CTA的執行後，新的CTA會被分配給它，直到所有待執行的CTA都執行完畢

［4］

。

一個SM當中包含了4個（視具體型號而不同）

SIMT單元(single-instruction, multiple-thread unit)

。同一個CTA中的執行緒將會按照32個一組的方式被分配給不同的SIMT unit執行。也即一個SM中最多同時執行128個active thread。需要注意的是，每個SIMT單元中，包含的計算元件數量只有16個（單精浮點16個，雙精8個），而待執行的執行緒卻有32個，因此一次單精浮點運算將要分兩批完成。（這部分是我瞎bb，在實際測試中，單精浮點耗時穩定為4週期，雙精在CTA執行緒數大於64之後耗時非線性增長，具體原因請自行腦補）。除此之外，每個SIMT單元還包含一定數量的LD/ST單元，負責提供資料訪問。

此外，據資料

［5］

所述，自Volta之後，存在名為Independent Thread Scheduling的新機制。該機制為每個執行緒設定獨立的PC和Call stack，這使得每個warp的執行緒能夠完全獨立的進行排程。

這一新機制導致了原有的許多預設warp同步的"高效能"程式碼需要被重新review以免導致錯誤的執行結果。

（怪不得之前讀文件見到過warp同步的機制）

但是需要注意的是，

這並不意味著SIMT的本質有所改變！同一時間一個SIMT unit仍然只能執行同一條指令！

如上圖，在Volta之前的架構中，屬於一個warp的執行緒，若要實現diverge，是透過active mask標記活躍執行緒來實現的。且一個warp公用一個PC與Call stack，這意味著一個then執行完畢前，else中的指令是沒有被髮射的可能的（因為沒有PC來記錄執行位置）。該種情況下，如果A指令是一條訪存指令，會帶來極大的latency，也沒法透過切換活躍執行緒先執行指令X的方式來掩蓋，而只能透過觸發warp級別的排程來掩蓋延遲。

而在Volta之後的架構中，得益於獨立執行緒分配的PC和Call stack。如果A指令因為訪存而被阻塞，排程器將有機會切換到else分支繼續執行指令X，並在隨後切換回指令B執行。這將會有效提高計算部件的利用率。

注意到，在V&T架構中，Z指令的執行可能不再同步，這會造成效能的下降。透過在指令Z之前插入__syncwarp指令，能夠強制指令Z同步執行，獲得性能提升。這是因為在多執行緒環境下，並不是任何時候擅自進行這樣的編譯最佳化都是合法的。

［6］

不管怎麼說，V&T架構相較於前，其分支效能多少還是有所提高。

而在實際的測試中發現，以上來自NV文件的內容真實（是）詳（狗）細（屁），

你爸爸永遠只告訴你小孩子才需要知道的

。

下面貼上實驗資料：

以下測試均在NV RTX 2080Ti上完成，編譯指令為：

nvcc -o test。exe test。cu -G -keep -arch compute_75 -code sm_75

測試用核心函式為：

__device__ float result；

__global__ void kernel_run2（float a， float b）{

float res = 1；

for （int i = 0； i < 100000； i++） {

res += i * res + a * a + 2 * b + b * b；

}

__syncthreads（）； //保留或去除這一行，對測試結果無影響。加這個只是因為考慮到可能是因為沒有使用任何的同步機制，所以仍有被悄咪咪最佳化的餘地？

result += res / 10000；

}

對應Turing架構生成程式碼，禁用所有最佳化。已檢查彙編確認運算量實打實的存在。且在執行實際測試部分前，已經透過預先執行一段gpu程式來warm up，使GPU在整個測試期間保持穩定的頻率和溫度。

按照已知的說法，當一個CTA中包含的執行緒數量增加時，每增加128個執行緒（4*WARP_SIZE or even 5*WARP_SIZE。 Whatever。），執行時間應當提高一倍。根據相關文章

［5］

，在Volta架構上，這一倍率的實際值為1。57。

而測試的結果很出乎意料：

每根曲線代表了不同的執行緒組數量下的結果，橫軸表示每個執行緒組包含的執行緒數，縱軸表示相對耗時。相對耗時已做舍入取整處理。

這完全可以用捉摸不透來形容，因為在期望中，不管執行緒組數量是多少，執行耗時應當隨執行緒組中執行緒數/128呈線性增長。

而該測試結果表現出兩個特點：

線上程組數量不多時，耗時完全不隨warp數量增加而增加。

線上程組數量足夠多時，耗時隨warp數量增加增長，但不表現為線性。

這似乎意味著文件中的說法與目前硬體對CTA的排程策略有著極大的出入。

特點2可以用CTA與程式設計模型中的執行緒組執行緒數量不一致來解釋。

透過合併數量太少的執行緒組至同一個CTA，並使用特殊的subCTA的同步機制（事實上根據文件描述的確有這樣的機制），仍可以保證在文件所述的硬體架構下，對程式設計師表現出合乎正常行為的程式設計模型。實際上，這樣的最佳化的確可以在程式設計師選擇的執行緒組中執行緒數量會導致極低的效能時，仍能透過不可見的最佳化保持較好的實際執行效能。這可能是的確存在於非開源部分的最佳化行為。

但是，特點1完全無法與文件表述相容。

因為CTA是執行緒通訊允許的最大範圍。而執行緒組是程式設計模型中允許的執行緒通訊的最大範圍。這意味著CTA只能比執行緒組大，而不可能比執行緒組小。而根據文件描述，一個CTA中的所有執行緒只能在同一個SM中執行。這意味著按文件說法，不論什麼情況下，增大執行緒組中執行緒數量必然會導致執行耗時增加。這是無法透過對程式設計師不可見的最佳化來避免的。

具體為什麼，天知道。這TM我報告沒法寫了。或者有哪位聚聚指出這份測試結果存在的問題，感激不盡！

猜測可能在Turing架構中，SM的概念已經淪為程式設計模型上的概念了？所有的執行緒其實都是統一排程的。故而在佔用滿所有的可用的SIMT Unit之前，耗時都不會增加？可是這麼大的改動怎麼會隻字未提呢。

下篇：

參考

^

Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking

^

register reuse cache

https：//github。com/NervanaSystems/maxas/wiki/SGEMM#calculating-c-register-banks-and-reuse

^

Turing instruction set

https：//docs。nvidia。com/cuda/cuda-binary-utilities/#turing

^

PTX Machine Model

https：//docs。nvidia。com/cuda/parallel-thread-execution/index。html#ptx-machine-model

^

a

b

Independent Thread Scheduling

https：//docs。nvidia。com/cuda/parallel-thread-execution/index。html#independent-thread-scheduling

^

volta architecture whitepaper