GPU編程和流式多處理器
流式多處理器（SM）是運行CUDA內核的GPU的一部分。本章重點介紹SM的指令集功能。
流式多處理器（SM）是運行我們的CUDA內核的GPU的一部分。每個SM包含以下內容。
? 可以在執行線程之間劃分的數千個寄存器
? 幾個緩存：
? –共享內存，用于線程之間的快速數據交換
? –恒定高速緩存，用于快速廣播恒定內存中的讀取
? –紋理緩存，以聚合紋理內存的帶寬
? – L1緩存，可減少對本地或全局內存的延遲
? Warp調度程序可以快速在線程之間切換上下文，并向準備執行的Warp發出指令
? 用于整數和浮點運算的執行核心：
? –整數和單精度浮點運算
? –雙精度浮點
? –用于單精度浮點先驗功能的特殊功能單元（SFU）
存在許多寄存器，以及硬件可以如此高效地在線程之間進行上下文切換的原因，可以最大程度地提高硬件的吞吐量。GPU被設計為具有足夠的狀態，以覆蓋執行等待時間和數百個時鐘周期的存儲等待時間，執行讀取指令后，設備存儲器中的數據可能需要數百個時鐘周期才能到達。
SM是通用處理器，設計與CPU中的執行內核有很大不同：目標時鐘頻率低得多；支持指令級并行性，但不支持分支預測或推測性執行；如果根本沒有緩存，則緩存較少。對于合適的工作負載，GPU中強大的計算能力足以彌補這些缺點。
自2006年推出首款支持CUDA的硬件以來，SM的設計一直在迅速發展，其代號為Tesla，Fermi和Kepler的三個主要修訂版本。開發人員可以通過調用cudaGetDeviceProperties（）并檢查cudaDeviceProp.major和cudaDeviceProp.minor或通過調用驅動程序API函數cuDeviceComputeCapability（）來查詢計算能力。計算能力1.x，2.x和3.x分別對應于Tesla類，Fermi類和Kepler類硬件。表8.1總結了每一代SM硬件中添加的功能。
表1 SM功能

以后會顯示不同SM的框圖。CUDA內核可以執行整數和單精度浮點指令。一個雙精度單元實現雙精度支持（如果可用）；和特殊功能單元實現倒數，倒數平方根，正弦/余弦和對數/指數函數。當執行指令所需的資源可用時，warp調度程序會將指令調度到這些執行單元。
本文重點介紹SM的指令集功能。因此，有時會引用“ SASS”指令，即ptxas或CUDA驅動程序將中間PTX代碼轉換成的本機指令。開發人員無法直接編寫SASS代碼；相反，NVIDIA通過cuobjdump實用程序使這些說明對開發人員可見，可以通過檢查編譯的微代碼來指導其源代碼的優化。

1. 存儲Memory
   1.1. 寄存器
   每個SM包含數千個32位寄存器，這些內核在啟動內核時指定給線程分配。寄存器是SM中最快，最豐富的內存。例如，開普勒類（SM 3.0）SMX包含65,536個寄存器或256K，而紋理緩存僅為48K。
   CUDA寄存器可以包含整數或浮點數據。對于能夠執行雙精度算術（SM 1.3和更高版本）的硬件，操作數包含在偶數值寄存器對中。在SM 2.0和更高版本的硬件上，寄存器對也可以保存64位地址。
   CUDA硬件還支持更廣泛的內存事務：內置int2 / float2和int4 / float4數據類型分別位于對齊的寄存器對或四元組中，可使用單個64位或128位寬的加載或存儲來讀取或寫入。一旦進入寄存器，就可以將各個數據元素稱為.x / .y（對于int2 / float2）或.x / .y / .z / .w（對于int4 / float4）。
   通過指定命令行選項–ptxas-options --verbose，開發人員可以使nvcc報告內核使用的寄存器數。內核使用的寄存器數量會影響SM中可容納的線程數量，通常必須仔細調整，以獲得最佳性能。可以使用–ptxas-options --maxregcount N指定用于編譯的最大寄存器數。
   注冊采樣Register Aliasing
   寄存器可以保存浮點或整數數據，某些內在函數僅用于迫使編譯器更改其變量視圖。__int_as_float（）和__float_as_int（）內部函數，導致32位整數和單精度浮點之間“改變性能”的變量。
   float__int_as_float（int i）;
   int __float_as_int（float f）;
   __double2loint（） ，__double2hiint（） ，和__hiloint2double（）內部函數類似原因寄存器來改變個性（通常就地）。__double_as_longlong（）和__longlong_as_double（）就地強制寄存器對；__double2loint（）和__double2hiint（）分別返回輸入操作數的最低和最高32位；__hiloint2double（）從高半部分和低半部分中構造一個雙精度型。
   int double2loint（double d）;
   int double2hiint（double d）;
   int hiloint2double（int hi，int lo）;
   double long_as_double（long long int i）;
   long long int __double_as_longlong（double d）;
   1.2. 本地Local Memory
   本地存儲器用于溢出寄存器，還用于存儲已索引且無法在編譯時計算其索引的局部變量。本地內存由與全局內存相同的設備內存池支持，因此它具有與Fermi和更高版本的硬件上的L1和L2緩存層次結構相同的延遲特性和優勢。本地內存的尋址方式是自動合并內存事務。硬件包括加載和存儲本地內存的特殊說明：SASS變體是Tesla的LLD / LST和Fermi和Kepler的LDL / STL。
   1.3. 全局Global Memory
   SM可以使用GLD / GST指令（在Tesla上）和LD / ST指令（在Fermi和Kepler上）讀取或寫入全局內存。開發人員可以使用標準的C運算符來計算和取消引用地址，包括指針算法和取消引用運算符*，[]和->。對64位或128位內置數據類型（int2 / float2 / int4 / float4）進行操作，自動使編譯器發出64位或128位加載和存儲指令。通過合并內存事務可實現最大的內存性能。
   特斯拉級硬件（SM 1.x）使用特殊的地址寄存器來保存指針。后來的硬件實現了一種加載/存儲架構，該架構使用相同的寄存器文件來存儲指針。整數和浮點值；以及用于恒定內存，共享內存和全局內存的相同地址空間。1個
   費米級硬件包括舊硬件不具備的一些功能。
   ? 通過“寬”加載/存儲指令支持64位尋址，其中地址保存在偶數寄存器對中。在32位主機平臺上不支持64位尋址。在64位主機平臺上，將自動啟用64位尋址。結果，針對為32位和64位主機平臺編譯的，相同內核生成的代碼，可能具有不同的寄存器計數和性能。
   ? L1緩存的大小可以配置為16K或48K。2（Kepler添加了將緩存拆分為32K L1 / 32K共享功能。）加載指令可以包括可緩存性提示（告訴硬件將讀取的內容拖入L1或繞過L1，并將數據僅保留在L2中）。可通過嵌入式PTX或通過命令行選項–X ptxas –dlcm = ca（默認設置在L1和L2中緩存）或–X ptxas –dlcm = cg（僅在L2中緩存）訪問這些文件。
   即使多個GPU線程在同一內存位置上運行，原子操作（或僅僅是“原子”）也可以正常工作地更新內存位置。在操作期間，硬件會在內存位置強制執行互斥。由于不能保證操作順序，因此通常支持的運算符是關聯的。3
   Atomics首先可用于SM 1.1和更高版本的全局內存以及SM 1.2和更高版本的共享內存。但是，在開普勒一代硬件出現之前，全局內存原子太慢而無法使用。
   當通過–gpu-architecture為nvcc指定了適當的體系結構時，表2中匯總的全局原子內在函數將自動變為可用。所有這些內在函數都可以對32位整數進行操作。SM 1.2中添加了對atomicAdd（），atomicExch（）和atomicCAS（）的64位支持。在SM 2.0中添加了32位浮點值（float）的atomicAdd（）。在SM 3.5中添加了對atomicMin（），atomicMax（），atomicAnd（），atomicOr（）和atomicXor（）的64位支持。
   表2原子操作
   

注意
由于原子操作是使用GPU的集成內存控制器中的硬件實現的，無法在PCI Express總線上運行，無法在與主機內存或對等內存相對應的設備內存指針上正確運行。
在硬件級別，原子有兩種形式：原子操作返回執行算子之前在指定內存位置的值，以及歸約操作，開發人員可以在內存位置“觸發并忘記”而忽略返回值值。由于如果不需要返回舊值，則硬件可以更有效地執行操作，編譯器將檢測是否使用了返回值，如果未使用，則發出不同的指令。例如，在SM 2.0中，指令分別稱為ATOM和RED。
1.4. 恒定內存Memory
常量內存駐留在設備內存中，但由另一個只讀緩存支持，該緩存經過優化，可以將讀取請求的結果廣播，到均引用同一內存位置的線程。每個SM均包含一個經過延遲優化的小型緩存，用于處理這些讀取請求。將內存（和緩存）設置為只讀可簡化緩存管理，因為硬件無需實施回寫策略來處理已更新的內存。
SM 2.x和后續硬件包括針對內存的特殊優化，該優化未表示為常量，但編譯器已將其標識為（1）只讀和（2）其地址不依賴于塊或線程ID。“均勻加載”（LDU）指令使用恒定的緩存層次結構讀取內存，并將數據廣播到線程。
1.5. 共享內存
共享內存非常快，是SM中的片上內存，線程可以將其用于線程塊內的數據交換。由于它是每個SM資源，因此共享內存的使用會影響占用率，即SM可以保留的warp數量。SM使用特殊指令加載和存儲共享內存：SM 1.x上為G2R / R2G，SM 2.x及更高版本上為LDS / STS。
共享內存安排為交錯的存儲體，通常針對32位訪問進行了優化。如果經線中有多個線程引用同一存儲體，則會發生存儲體沖突，并且硬件必須連續處理內存請求，直到為所有請求提供服務為止。為了避免存儲區沖突，應用程序將基于線程ID以交錯模式訪問共享內存，如下所示。
extern shared float shared[];
float data = shared[BaseIndex + threadIdx.x];
從相同的32位共享內存位置讀取扭曲中的所有線程也很快。硬件包括廣播機制以針對這種情況進行優化。寫入同一存儲區的操作會由硬件進行序列化，降低性能。寫入同一地址會導致爭用情況，應避免。
對于2D訪問模式（例如圖像處理內核中的像素圖塊），最好填充共享內存分配，以便內核可以引用相鄰行，而不會引起存儲體沖突。SM 2.x和后續硬件具有32個存儲區，其中每個存儲區可供2個圖塊使用，其中同一warp中的線程可以按行訪問數據，這是將圖塊大小填充為33個32位字的倍數的好策略。
在SM 1.x硬件上，共享內存的大小約為16K。在更高版本的硬件上，總共有64K的L1高速緩存，可以配置為16K或48K共享內存，其余部分用作L1高速緩存。
在過去的幾代硬件中，NVIDIA改進了硬件對除32位以外的操作數大小的處理。在SM 1.x硬件上，來自同一存儲區的8位和16位讀取導致存儲區沖突，而SM 2.x和更高版本的硬件可以從同一存儲區廣播中，任何大小的讀取。同樣，共享內存中的64位操作數（例如double）比SM 1.x上的32位操作數慢得多，以至于開發人員有時不得不求助于將數據分別存儲為上下半部分。SM 3.x硬件為主要在共享內存中使用64位操作數的內核添加了一項新功能：將存儲體大小增加到64位的模式。
共享內存中的原子算子
SM 1.2添加了在共享內存中執行原子操作的功能。全局內存與使用單個指令（取決于GATOM或GRED，取決于是否使用返回值）來實現原子的全局存儲器不同，共享內存原子是通過顯式的鎖定/解鎖語義實現的，編譯器發出的代碼會導致每個線程循環這些鎖定操作，直到線程執行其原子操作為止。
清單1將源代碼提供給atomic32Shared.cu，該程序專門用于編譯以突出顯示共享內存原子的代碼生成。清單2顯示了為SM 2.0生成的最終微代碼。注意，LDSLK（與鎖共享負載）指令如何返回謂詞，該謂詞說明是否獲取了鎖，確定了要執行更新的代碼，并且代碼循環運行，直到獲取了鎖并執行了更新。
鎖是按32位字執行的，鎖的索引由共享內存地址的位2–9確定。注意避免爭用，否則清單2中的循環最多可以迭代32次。
list1. atomic32Shared.cu
global void
Return32( int *sum, int *out, const int *pIn )
{
extern shared int s[];
s[threadIdx.x] = pIn[threadIdx.x];
__syncthreads();
(void) atomicAdd( &s[threadIdx.x], *pIn );
__syncthreads();
out[threadIdx.x] = s[threadIdx.x];
}
Listing 2. atomic32Shared.cubin (microcode compiled for SM 2.0)
code for sm_20
Function : _Z8Return32PiS_PKi
/0000/ MOV R1, c [0x1] [0x100];
/0008/ S2R R0, SR_Tid_X;
/0010/ SHL R3, R0, 0x2;
/0018/ MOV R0, c [0x0] [0x28];
/0020/ IADD R2, R3, c [0x0] [0x28];
/0028/ IMAD.U32.U32 RZ, R0, R1, RZ;
/0030/ LD R2, [R2];
/0038/ STS [R3], R2;
/0040/ SSY 0x80;
/0048/ BAR.RED.POPC RZ, RZ;
/0050/ LD R0, [R0];
/0058/ LDSLK P0, R2, [R3];
/0060/ @P0 IADD R2, R2, R0;
/0068/ @P0 STSUL [R3], R2;
/0070/ @!P0 BRA 0x58;
/0078/ NOP.S CC.T;
/0080/ BAR.RED.POPC RZ, RZ;
/0088/ LDS R0, [R3];
/0090/ IADD R2, R3, c [0x0] [0x24];
/0098/ ST [R2], R0;
/00a0/ EXIT;
…
1.6. 障礙和連貫性
熟悉的__syncthreads（）內部函數會等待，直到繼續執行線程塊中的所有線程為止。需要保持線程塊內共享內存的一致性。其他類似的內存屏障指令，也可以用于在更大范圍的內存上執行某些排序，如表3所述。
表3內存屏障本質

總結

以上是生活随笔為你收集整理的GPU编程和流式多处理器的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。