大家做高性能計算的朋友，想必對CPU的執行模式已經非常熟悉了吧。當代高級些的CPU一般采用超標量流水線，使得毗鄰幾條相互獨立的指令能夠並行執行——這稱為指令集並行（ILP，Instruction-Level Parallelism）；而像x86引入的SSE（Streaming SIMD Extension）、AVX（Advanced Vector Extension），以及ARM的NEON技術都屬於數據級並行（Data-Level Parallelism）。而GPGPU的執行與CPU比起來還是有不少差異的。這里，為了能夠讓大家更好地理解、並使用OpenCL，想談談當前主流用於超算的GPGPU的執行模式。
下面主要針對nVidia的Fermi和Kepler架構以及AMD的TeraScale3（Radeon HD 6900系列）和GCN架構進行分析。
我們先來簡單介紹一下OpenCL中的一些術語對應到nVidia以及AMD的GPGPU硬件中的稱謂。
在物理上，一個GPGPU作為一個計算設備，在OpenCL中就稱為Device。在這么一個計算設備中由若干大的計算核心構成，這個大的計算核心在OpenCL中稱為CU（Compute Unit），在nVidia中稱為SM（Streaming Multiprocessor），在AMD中則稱為SIMD（Single-Instruction-Multiple-Data）。一個大的計算核心中又由許多小的計算核心構成，這種小的計算核心在OpenCL中稱為PE（Processing Element），在nVidia和AMD中均稱為SP（Stream Processor）。此外，nVidia也好、AMD也罷，甚至還有一些如Intel HD Graphics這樣的GPGPU，有一個OpenCL中沒有對應到的術語，它其實屬於一種GPGPU的執行模式，我這里暫且稱為“線程流”。以應用開發者的角度來看，它是GPGPU中線程執行的最小並行粒度，稍后會詳細講解。這個概念在nVidia中稱為Warp，在AMD中稱為Wavefront。
我們下面先來看看nVidia的Fermi架構是如何在OpenCL中執行的。在Fermi架構中，一個SM一共有32個SP，16個存儲器讀寫單元，四個特殊功能計算單元SFU（用於計算超越函數等復雜操作），64KB的共享存儲器（Local Memory），32768個32位寄存器，兩個Warp調度器與兩個指令分派單元。
在此架構中，SM調度線程時會將32個線程（OpenCL中稱為work-item）組成一組，然后並發執行。這個32線程組就是一個warp。由於每個SM含有兩個warp調度器與兩個指令分派單元，因此這就能夠將兩個warp同時進行發射和執行。Fermi的雙warp調度器先選擇兩個warp，然后從每個warp發射一條指令到一個十六核心的組，或是十六個讀寫單元，或是四個SFU。正由於warp執行是獨立的，因此Fermi的調度器無需檢查指令流內的依賴性。
那么這個執行模式如何映射到一個OpenCL的kernel程序里呢？我們現在假設給kernel分配了512個work-item（Fermi架構的Max Work-group Size為1024），work-group size也是512，然后執行以下kernel代碼：
1 __kernel ocl_test(__global int *p) 2 { 3 int index = get_global_id(0); 4 5 int x = p[index]; 6 7 x += 10; 8 9 p[index] = x; 10 }
我們首先要知道，一個work-group是被一個SM負責執行的。因為一個work-group中所含有的寄存器以及local memory資源都屬於一個SM里的資源。所以，對於上述配置，這512個work-item都將在一個SM中完成執行。
那么上面提到的warp調度器從每個warp發射一條指令到一個十六核心的組，或是十六個讀寫單元，是怎么回事呢？
我們之前提到，一個SM一共有32個核心，每個調度器會將指令發送到其中一半（相應的16個核心），這樣兩個調度器同時發射一次，那么相應的指令正好能在這32個核心中執行一次。因此，對於warp調度器而言，完整地執行一個warp會將一條指令連續發射兩次。而這兩次發射對於程序員而言可以看作是原子的，即不可分割的。對於每個調度器，前一次發射，執行前16個work-item；后一次發射，執行后16個work-item；這前后兩組16個work-item就組成了一個完整的32個work-item的Warp。
下面，我們看上述kernel程序，第5行、第9行，核心執行的是讀寫操作；第7行核心執行的算術計算。那么global ID則是從0到511。如果將一個warp的一次執行看作為一個周期的話，那么：
第一個周期：id從0－31的work-item，組成為warp0；id從32到63的work-item組成為warp1，送到SM同時執行一次。
第二個周期：id從64－95的work-item，組成為warp2；id從96到127的work-item組成為warp3，送到SM同時執行一次。
...
第八個周期：id從448－479的work-item，組成為warp14；id從480到511的work-item組成為warp15，送到SM同時執行一次。
這樣，8個周期就將整個work-group執行了一遍。
Kepler架構的執行跟Fermi差不多，不過原本的SM，現在更名為SMX，它擁有四個warp調度器和八個指令分派單元，這就允許warp調度器選擇4個warp被並發發射執行，而且又因為每個warp調度器又對應兩個指令分派單元，從而使得每個warp的鄰近兩條相互獨立的指令能夠在一個周期內被同時執行。一個SMX至少含有128個核心，因此能夠並行執行的work-item數量都是Fermi架構的兩倍。因此，對於Fermi架構而言，我們在分配一個work-group size的時候，應該將它分配為64的倍數；而Kepler架構，則應該是128的倍數，這樣能充分利用調度器而達到峰值計算。
下面我們再來看看Radeon HD Graphics TeraScale3的執行方式。Radeon HD Graphics 6900系列的一個SIMD作為一個CU。每個CU含有16個SP，256KB的寄存器文件（65536個32位寄存器），32KB的Local Memory。其中，每個SP含有四個獨立的算術邏輯單元（ALU），允許四條相互獨立的標量數據計算同時執行。不過每個SIMD的線程調度器僅一個。
Radeon HD Graphics的執行是以wavefront的模式執行的。一個SP中的每個ALU對應一條獨立的wave，這樣，一個SIMD中的每個SP就可以在一個周期以四條wave同時執行，當然前提是這四個操作相互獨立，並且正好能被編排到一個SP中的各個ALU中。一個周期同時執行16個SP，這樣最多就能完成64個標量算術邏輯操作。而GPU對同一組指令連續發射4次就正好把一條完整的wavefront執行完成，一共最多能完成16條wave，256個標量算術邏輯操作。下面舉一個OpenCL的kernel例子：
假定，設置work-group大小為256，一共256個work-item。
1 __kernel ocl_test(__global int *a, __global int *b) 2 { 3 int index = get_global_id(0); 4 5 int4 vecA = vload4(index, a); 6 int4 vecB = vload4(index, b); 7 8 vecA.x += vecB.x 9 vecA.y += vecB.y; 10 vecA.z += vecB.z; 11 vecA.w += vecB.w; 12 }
上述代碼第8到第11行，我們是把vecA += vecB;這條向量計算語句拆成了四條標量語句。OpenCL驅動確實也是如此做的，這樣，這四條獨立的標量算術操作正好能對應上一個SP的四個ALU上。就拿這四條語句而言：
第一個周期：id為0－15的work-item，每個work-item的第8行對應wave0，第9行對應wave1，第10行對應wave2，第11行對應wave3；
第二個周期：id為16－31的work-item，每個work-item的第8行對應wave4，第9行對應wave5，第10行對應wave6，第11行對應wave7；
第三個周期：id為32－47的work-item，每個work-item的第8行對應wave8，第9行對應wave9，第10行對應wave10，第11行對應wave11；
第四個周期：id為58－63的work-item，每個work-item的第8行對應wave12，第9行對應wave13，第10行對應wave14，第11行對應wave15；
這樣，這四個周期完整地執行了整條wavefront，一共占用64個work-item，執行了256次算術操作。當然，對於應用開發者而言，這四個周期是原子的，不可被分割的。這也是為啥這四個周期執行了獨立的16條wave的緣由。所以，對於VLIW4或VLIW5架構的Radeon HD Graphics，我們設置work-group size最好是64的倍數。
第五個周期：id為64－79的work-item，每個work-item的第8行對應wave0，第9行對應wave1，第10行對應wave2，第11行對應wave3；
...
第十六個周期：id為240－255的work-item，每個work-item的第8行對應wave12，第9行對應wave13，第10行對應wave14，第11行對應wave15；
這樣，16個周期就完成了所有256個work-item對第8到第11行語句的執行，總共執行了1024次算術操作。當然，這是在最好的情況下。倘若第8到第11行有些語句存在相互依賴，那么將會導致某些操作無法被同時放入SP的四個ALU單元，從而使得SP在執行時某些ALU計算單元處於空閑狀態。這也是為啥以VLIW類型進行執行的GPGPU，在寫OpenCL代碼時最好使用向量數據類型進行操作的原因。通常，一個向量的每條通道（lane）相互獨立，使得它們能夠被送到SP的各個ALU中。
下面接着談一下AMD的GCN架構。GCN架構與TeraScale系列完全不同，反而跟Kepler架構更接近。GCN架構將CU這個概念正式運用到了GPU的硬件架構上。原本，TeraScale3的單個SIMD由四個獨立的ALU組成，而被改成了GCN中，一個CU由四個SIMD單元組成，其中，每個SIMD僅由單個ALU構成。每個SIMD還含有獨立的64KB的寄存器，整個CU含有64KB的Local Memory。
在GCN架構中，一個SIMD至少對應10條wavefront，那么對於一個CU而言就是40條wavefront（4個SIMD，每個SIMD有10條wavefront）可以在執行流水線上運行。而每條wavefront對應執行64個work-item，並且可以在各自不同的work-group上，甚至不同的kernel上執行。那么一個CU則一次可執行2560個work-item。而每個SIMD一次可同時執行16個work-item，而且每個SIMD可以對各自的wavefront進行操作。這樣，GCN的執行模式與Kepler就很像了。
在GCN架構中，指令分發序列器以每個CU為單位進行分發，它管理4階段的執行。也就是說，一個SIMD在執行完一整條wavefront與之前的VLIW4一樣，需要連續發射4次完成。
其中，4個CU組成一個簇共享一個32KB的四路組相聯的L1指令Cache，通過L2 Cache進行后備緩存。Cache行為64字節長，一般能保留8條指令。當Cache滿的時候，系統會發出一條新的請求，以最近最少使用策略（LRU）將某條Cache行逐出，為新的指令留出空間。4個CU所共享的L1 Cache含有4個段，並且可以維持每周期對所有4個CU做取32字節指令操作（每個CU取一條8字節的指令）。去指令在一個CU內的4個SIMD之間進行仲裁，基於工作時長、調度優先級以及對wavefront指令緩存的利用。
一旦指令取到wavefront緩存中，下一步就進行譯碼並發射指令。CU在每個周期，通過輪詢仲裁方式選擇一個SIMD來譯碼並發射。所選中的SIMD可以從10條wavefront的緩存中譯碼並發射多達5條指令到執行單元。此外，在wavefront緩存中還可以執行一條特殊功能指令（比如，NOP操作、柵欄操作、暫停操作、跳過一條謂詞向量指令等），而不占用任一功能單元。每個CU具有16個緩存來追蹤柵欄指令。柵欄指令會迫使一條wavefront進行全局同步。
CU前端可以譯碼並發射七種不同類型的指令：分支，標量ALU或訪存，向量ALU，向量訪存，LDS（Local Data Share，相當於OpenCL中的Local Memory）訪問，全局數據訪存，特殊功能指令。每個SIMD每個周期只能發射每種類型的其中一條指令，以避免過多注冊（oversubscribing）執行流水線。為了維護順序執行，每條指令也必須來自不同的wavefront。每個SIMD具有10條wave，那么一般就有許多種選擇了。除了上述這兩種限制，任意混合都是被允許的，這給了編譯器充分的自由來安排指令發射執行。
CU的前端可以每個周期發射5條指令到一個6個向量與標量相混合的執行流水線，使用兩個寄存器文件。向量單元提供了對於圖形着色器以及計算密集的通用目的應用強大的計算能力。兩個標量單元與指令緩存中處理的特殊指令一起負責GCN架構所有的控制流。
每個CU含有8KB的標量寄存器文件，給每個SIMD划分為512個條目。在一個SIMD上，所有10條wavefront共享這些標量寄存器。一條wavefront可以分配112個用戶寄存器以及若干保留作為架構狀態的寄存器。每個寄存器是32位寬，並且鄰近兩個寄存器可以用於存放一個64位的值。
對於向量寄存器，由於每個SIMD獨立地執行一條wavefront，因此一個CU中的寄存器文件可以被划分為四個獨立的片段。
向量通用目的寄存器（vGPR）包含了64個通道（lane），每個通道寬度為32位。鄰近的vGPR可以被聯結為64位或128位數據。每個SIMD具有vGPR的64KB子部分，一個CU所占用的向量寄存器總數是固定的。每個SIMD的子部分被細粒度地分段，並且可以同時讀X寄存器，寫Y寄存器。
每個SIMD包含了一條16通道（lane）的向量流水線。每條通道可以執行一個單精度的融合或非融合乘加操作或是一個24位整數操作。一條wavefront在一單個周期被發送到一個SIMD，不過要花費4個周期來執行所有64個work-item的的執行操作。
同時，我們之前已經提到了，“在GCN架構中，指令分發序列器以每個CU為單位進行分發，它管理4階段的執行。每個SIMD一次可同時執行16個work-item，而且每個SIMD可以對各自的wavefront進行操作。一個CU在每個周期，通過輪詢仲裁方式選擇一個SIMD來譯碼並發射”。因此，在一個CU內，四個SIMD是以4級流水線那樣被調度執行的。當然，這個調度不是嚴格按照某一次序，作為一個例子，我們可以想象第一個周期把第一條指令發射給SIMD0執行；第二個周期把第1條指令發射給SIMD1執行，同時SIMD0執行第2條指令；第三個周期，把第1條指令發射給SIMD2執行，同時SIMD0執行第3條指令，SIMD1執行第2條指令；第四個周期，把第一條指令發射給SIMD3執行，同時，SIMD0執行第4條指令，SIMD1執行第3條指令，SIMD2執行第2條指令。這樣，當整個流水線被填滿時候，該CU即處於峰值計算狀態，四個周期即能同時對一個CU的4個SIMD同時發射執行，完成4整條完整的wavefront。這里還要注意的是，同一種指令被分配到不同的wavefront執行時是屬於不同的指令。其實，每個CU中含有4個獨立的指令緩存，所以對於每個SIMD正好可以使用一個。就拿上述代碼第8條指令“vecA.x += vecB.x”而言，這一條指令可以在發射到一個CU的不同SIMD執行時，其實是被復制了4份，分別放入到該CU的4個獨立的指令緩存中。可見，GCN架構的GPU在執行模式上顯得十分靈活。
因此，我們對於GCN架構的GPGPU，我們可以把它想象成一個通用的CPU。每個CU如果看成一個核心的話，那么其中的40條wavefront可以被看作為40個硬件線程（類似於HTT，超線程技術），通過4級流水線執行；而每條wavefront又是以SIMD的方式執行的，4個周期能處理完一條wavefront，共64個work-item。而我們給每條wavefront發射的指令當然都是相互獨立且不同的，盡管可能都是同一種，比如上面的“vecA.x += vecB.x”。而在一條wavefront內部，每個work-item所執行的指令絕對是同一條，僅僅是數據通道不同（64條lane）。
參考資料：