李曉雯　崔翔　殷瑞杰　劉強

摘 要： 討論在Fermi結構GPU使用CUDA對GEMM（單精度和雙精度）算法進行優(yōu)化，以及Fermi體系結構的新特性（如緩存）對性能的影響。GPU緩存一方面可以提高處理器在運行時數據訪問的局部性，另一方面使得代碼性能對與性能相關算法參數的依賴變得不可預測。自動優(yōu)化技術可以用來解決這一問題。自動優(yōu)化的SGEMM和DGEMM代碼在Tesla C2050 GPU上達到了563GFlops和253GFlops的性能。代碼使用CUDA和C語言進行實現，未進行二進制代碼級別的優(yōu)化。

關鍵詞： GPU程序設計； 矩陣乘法； 自動優(yōu)化； GEMM模板

中圖分類號： TN40； TP312 文獻標識碼： A 文章編號： 1004?373X（2014）10?0137?04

Abstract： Automatic optimization of GEMM algorithm on Fermi GPU and the impact of Fermi GPUs architectural features on performance are discussed in this paper. The cache on GPU can not only improve the data access locality of processors， but also make the code performance relying on relative algorithm parameters unpredictable. Auto?tuning can be used to solve this problem. Auto?tuned SGEMM and DGEMM codes achieve 563 GFlops and 253 GFlops respectively on C2050 GPU. The codes are implemented by using CUDA and C language， but the optimization on the binary level is not involved.

Keywords： GPU programming； matrix multiplication； automatic optimization； GEMM template

0 引 言

Fermi是Nvidia公司支持CUDA編程模型的新一代GPU。與GT200體系結構相比，新型的Tesla 2050 GPU具有一些新的特性：如增強的雙精度浮點性能、L1/L2 緩存結構、更多的寄存器、更大的共享存儲器、ECC支持和更快的原子操作[1?3]。

由于Tesla 2050和GT200使用同樣的編程模型，因此程序員期望在GT200上得到良好性能優(yōu)化的代碼也一樣可以在Tesla 2050上取得良好的性能。實際上，程序員依然需要調整他們運行在GT200上的代碼以在Tesla 2050上取得最高的性能。在Tesla 2050上，雖然每一個MP的寄存器文件被加倍了，但由于每一個MP的核數由8增加為32，因此實際上每個線程可用的寄存器數目實際上減少了一半。這使得程序員需要更加注意寄存器的使用效率。新增加的緩存結構一方面帶來了運行時數據訪問局部性的好處，另一方面也增加了代碼性能的不可預測性。程序員依然需要了解GPU的硬件特性以得到高效的GPU代碼。

自動調節(jié)技術是在復雜和不可預測的體系結構上得到近似最優(yōu)代碼的一種實用的技術。使用自動調節(jié)技術得到的SGEMM和DGEMM代碼在Tesla 2050分別達到了563 GFlops和253 GFlops的速度，相對于CUBLAS 3.0分別具有1.7倍和1.6倍的加速[4]。

1 Fermi的新特性

1.1 L1/L2緩存

與GT200體系結構相比，Fermi增加了L1/L2緩存以提高訪問設備存儲器的性能，如圖1所示。程序員對L1緩存的使用可以進行控制：64 KB的片上存儲可以被用于L1緩存或共享存儲器，而16 KB或48 KB存儲分別用于L1緩存或共享存儲器（抑或相反）可以在每次內核調用時進行控制。使用到局部存儲器的內核代碼可以從新增的L1緩存中受益。除了L1緩存，Fermi還提供768 KB的L2緩存。CUDA程序設計模型的原有特點是暴露硬件體系結構使得程序員對代碼的性能可以進行良好的控制，而緩存的引入卻使得CUDA代碼的行為和性能變得難以預測。甚至CUDA編程手冊[3]都建議程序員通過實驗的方法來確定L1緩存或共享存儲器的配置問題。對于L2緩存，一個簡單的使得代碼受益的方法是保證訪問相同設備存儲器地址空間的線程塊被連續(xù)的調度；這可以通過將blockIdx.x和blockIdx.y變量進行對調實現。

考慮到緩存效果的不可預測性，自動優(yōu)化技術可以用來得到近似最優(yōu)的CUDA代碼。首先將算法實現的代碼進行參數模板化，通過選擇不同的參數組合來自動得到具有良好性能的代碼。

1.2 寄存器文件

在文獻[4?5]中提到，相對于共享存儲器而言，在設計算法時應該優(yōu)先選擇使用寄存器以得到良好的性能。在GT200體系結構上，使用寄存器間的MAD指令可以達到98%的理論運算性能。在Fermi發(fā)布之前，程序員都期望在新的GPU體系結構上，寄存器文件的大小能被增加，從而使得先前的代碼在Fermi上能夠取得更優(yōu)的性能。

與GT200體系結構相比，Fermi上每個多處理器的寄存器文件的大小由16 KB增加為32 KB；而與此同時，每個多處理器的微核數目從8增加到32。這意味著在Fermi體系結構上，每個多處理器的微核的可用寄存器實際上是減少了。在GT200體系結構上，每個多處理器需要256個活動線程以掩蓋指令流水線的延遲，而在Fermi體系結構上，則需要更多的活動線程來掩蓋指令流水線的延遲。

1.3 32/64位設備代碼

在Fermi體系結構上，如果代碼按照64?bit的模式編譯，則CUDA編譯器會將CPU代碼和設備代碼都編譯成為64?bit的目標代碼。在這種情況下，設備代碼中的指針變量將會占用多出一倍的寄存器空間。由于Tesla C2050的設備存儲器容量不超過4 GB（在增加ECC的情況下可用的設備存儲器空間只有2.625 GB），因此完全沒有必要在設備代碼中使用64 b的指針。因此，在本文實現的Fermi體系結構上的GEMM代碼中，CPU代碼和設備代碼總是被分別編譯的。

1.4 設備存儲器訪問

在GT200體系結構上，對設備存儲器的訪問是按照半個warp的單位來進行處理的，而在Fermi體系結構上，是按照一個warp的單位來進行處理的。因此，程序員需要調整內核調用時的維度設置。對于具有兩個維度的線程塊，其x維度的大小應該是warp大小的整倍數，而非半個warp大小的整倍數，從而使得每一個warp在訪問設備存儲器時可以得到較高的性能。

1.5 Bank沖突

在GT200體系結構上，共享存儲器具有16個bank，而且對其的訪問是按照半個warp的單位來進行處理的，而在Fermi體系結構上，共享存儲器具有32個bank，而且對其的訪問是按照整個warp的單位來進行處理的。每一個bank的大小為32位。

1.6 多內核并行執(zhí)行

Fermi體系結構支持多個內核代碼的并行執(zhí)行，使得不同應用上下文的內核代碼可以同時在一個GPU上運行；這樣，多個小的內核代碼可以共同利用一個GPU上的計算資源。這也是Fermi體系結構的新特性，但在本文自動優(yōu)化的GEMM代碼中并未使用到。

2 自動優(yōu)化的GEMM代碼

2.1 GEMM代碼模板

文獻[4]描述了在GTX280實現的達到393 Gflops性能的SGEMM內核代碼，在此，依然使用該代碼作為實現自動優(yōu)化的代碼模板。在該實現中，一個Csub被一個線程塊進行計算，根據線程塊中線程數目的多少，一個線程可以計算Csub的半列或多列元素。例如，對于m= 16和n=64，而線程塊具有16×4個線程，則每一個線程將計算一整列Csub的元素；如果線程塊具有16×8個線程，則每一個線程將計算半列Csub的元素。與文獻[4]中的實現類似，每一個線程在使用線程ID計算出訪問矩陣A、B和C的指針位置之后，進入一個循環(huán)。在每一輪循環(huán)中，一個線程塊從設備存儲器中讀入一個Asub的數據到共享存儲器中，之后，又一個內層的循環(huán)被執(zhí)行：在每一輪內層循環(huán)中，一個線程從矩陣B中讀入一個或多個元素，把這些元素與共享存儲器中相應的數據做計算，將結果累加到Csub對應的寄存器中。最后，每個線程將其計算的Csub的數據寫回到設備存儲器中。

盡管與文獻[4]中的GEMM實現具有類似的代碼結構，但是為了滿足緩存的友好性，對該代碼模板進行了一個重要的修正。在此計算中，矩陣B總是位于設備存儲器中，一列線程塊總是需要讀取一列Bsub的數據，如圖2（a）所示。如果線程塊能夠按照列優(yōu)先的順序調度，則可以達到更好的緩存命中效果。因此，在此GEMM模板中，blockIdx.x和blockIdx.y變量的順序被對調，從而達到如圖2（b）所示的設備存儲器訪問效果。這樣，在進行這個轉置之后，線程塊被以較優(yōu)的緩存命中效果的方式調度。在后面給出的性能測試中，將對比這個對調進行或不進行的性能結果。

2.2 對代碼模板進行自動優(yōu)化

設計的自動調節(jié)程序根據代碼模板生成代碼并測試其性能結果。在代碼模板中，5個參數（m，k，n，tx和ty）決定代碼模板的執(zhí)行行為；此外，引入2個額外的參數，一個用來決定線程塊的維度是否對調，來測試緩存的效果，另一個確定L1緩存和共享存儲器的配置比例。因此，整個代碼模板的行為是由7個參數確定的。

圖3顯示經過優(yōu)化得到的矩陣大小為2 048時的DGEMM代碼的實例。注意：在該代碼中，blockIdx.x和blockIdx.y變量通過C語言的宏進行了對調，而相應的調用代碼也做了相應的改變。在此代碼中，m=8，k= 64，n=1 024，tx=64，ty=8，因此一個線程負責計算Csub中的兩列數據。

3 性能測試

4 結 語

為了在Fermi體系結構上書寫高效的代碼，程序員需要很好地了解Fermi體系結構的新特性，以及它們是如何影響程序的性能的。對于Fermi體系結構，程序員尤其要關心緩存對性能的影響。對于Fermi這種復雜的行為和性能難以預測的硬件體系結構，自動優(yōu)化技術不失為得到高性能代碼的一種實用技術。

參考文獻

[1] NVIDIA Corp. Whitepaper： NVIDIA's next generation CUDA compute architecture [R/OL]. [2012?05?18]. http：//www. insidehpc.com.

[2] NVIDIA Corp. Tuning CUDA applications for Fermi [R/OL]. [2011?05?03]. http：// www. people.maths.ox.ac.uk/gilesm/cuda/doc/Fermi_Tuning_Guide.

[3] NVIDIA Corp. CUDA compute unified device architecture， programming guide， Version 3.0 [R/OL]. [2010?05?03]. http：// www. mohamedfahmed.wordpress.com

[4] CUI Xiang， CHEN Yi?feng， MEI Hong， et al. Auto?tuning GEMM for GPGPU with Cache [C]// Proceedings of 2010 IEEE 16th International Conference on Parallel and Distributed Systems （ICPADS）. Shanghai， China： IEEE， 2010： 237?242.

[5] VOLKOV V， DEMMEL J W. Benchmarking GPUs to tune dense linear algebra [C]// Proceedings of 2008. International Conference for High Performance Computing， Networking， Storage and Analysis Austin， TX： [s.n.]， 2008： 1?11.

[6] 李曉雯，崔翔.GPU矩陣乘法和FFT算法的性能優(yōu)化[J].現代電子技術，2013，36（4）：80?84.

1.3 32/64位設備代碼

在Fermi體系結構上，如果代碼按照64?bit的模式編譯，則CUDA編譯器會將CPU代碼和設備代碼都編譯成為64?bit的目標代碼。在這種情況下，設備代碼中的指針變量將會占用多出一倍的寄存器空間。由于Tesla C2050的設備存儲器容量不超過4 GB（在增加ECC的情況下可用的設備存儲器空間只有2.625 GB），因此完全沒有必要在設備代碼中使用64 b的指針。因此，在本文實現的Fermi體系結構上的GEMM代碼中，CPU代碼和設備代碼總是被分別編譯的。

1.4 設備存儲器訪問

在GT200體系結構上，對設備存儲器的訪問是按照半個warp的單位來進行處理的，而在Fermi體系結構上，是按照一個warp的單位來進行處理的。因此，程序員需要調整內核調用時的維度設置。對于具有兩個維度的線程塊，其x維度的大小應該是warp大小的整倍數，而非半個warp大小的整倍數，從而使得每一個warp在訪問設備存儲器時可以得到較高的性能。

1.5 Bank沖突

在GT200體系結構上，共享存儲器具有16個bank，而且對其的訪問是按照半個warp的單位來進行處理的，而在Fermi體系結構上，共享存儲器具有32個bank，而且對其的訪問是按照整個warp的單位來進行處理的。每一個bank的大小為32位。

1.6 多內核并行執(zhí)行

Fermi體系結構支持多個內核代碼的并行執(zhí)行，使得不同應用上下文的內核代碼可以同時在一個GPU上運行；這樣，多個小的內核代碼可以共同利用一個GPU上的計算資源。這也是Fermi體系結構的新特性，但在本文自動優(yōu)化的GEMM代碼中并未使用到。

2 自動優(yōu)化的GEMM代碼

2.1 GEMM代碼模板

文獻[4]描述了在GTX280實現的達到393 Gflops性能的SGEMM內核代碼，在此，依然使用該代碼作為實現自動優(yōu)化的代碼模板。在該實現中，一個Csub被一個線程塊進行計算，根據線程塊中線程數目的多少，一個線程可以計算Csub的半列或多列元素。例如，對于m= 16和n=64，而線程塊具有16×4個線程，則每一個線程將計算一整列Csub的元素；如果線程塊具有16×8個線程，則每一個線程將計算半列Csub的元素。與文獻[4]中的實現類似，每一個線程在使用線程ID計算出訪問矩陣A、B和C的指針位置之后，進入一個循環(huán)。在每一輪循環(huán)中，一個線程塊從設備存儲器中讀入一個Asub的數據到共享存儲器中，之后，又一個內層的循環(huán)被執(zhí)行：在每一輪內層循環(huán)中，一個線程從矩陣B中讀入一個或多個元素，把這些元素與共享存儲器中相應的數據做計算，將結果累加到Csub對應的寄存器中。最后，每個線程將其計算的Csub的數據寫回到設備存儲器中。

盡管與文獻[4]中的GEMM實現具有類似的代碼結構，但是為了滿足緩存的友好性，對該代碼模板進行了一個重要的修正。在此計算中，矩陣B總是位于設備存儲器中，一列線程塊總是需要讀取一列Bsub的數據，如圖2（a）所示。如果線程塊能夠按照列優(yōu)先的順序調度，則可以達到更好的緩存命中效果。因此，在此GEMM模板中，blockIdx.x和blockIdx.y變量的順序被對調，從而達到如圖2（b）所示的設備存儲器訪問效果。這樣，在進行這個轉置之后，線程塊被以較優(yōu)的緩存命中效果的方式調度。在后面給出的性能測試中，將對比這個對調進行或不進行的性能結果。

2.2 對代碼模板進行自動優(yōu)化

設計的自動調節(jié)程序根據代碼模板生成代碼并測試其性能結果。在代碼模板中，5個參數（m，k，n，tx和ty）決定代碼模板的執(zhí)行行為；此外，引入2個額外的參數，一個用來決定線程塊的維度是否對調，來測試緩存的效果，另一個確定L1緩存和共享存儲器的配置比例。因此，整個代碼模板的行為是由7個參數確定的。

圖3顯示經過優(yōu)化得到的矩陣大小為2 048時的DGEMM代碼的實例。注意：在該代碼中，blockIdx.x和blockIdx.y變量通過C語言的宏進行了對調，而相應的調用代碼也做了相應的改變。在此代碼中，m=8，k= 64，n=1 024，tx=64，ty=8，因此一個線程負責計算Csub中的兩列數據。

3 性能測試

4 結 語

為了在Fermi體系結構上書寫高效的代碼，程序員需要很好地了解Fermi體系結構的新特性，以及它們是如何影響程序的性能的。對于Fermi體系結構，程序員尤其要關心緩存對性能的影響。對于Fermi這種復雜的行為和性能難以預測的硬件體系結構，自動優(yōu)化技術不失為得到高性能代碼的一種實用技術。

參考文獻

[1] NVIDIA Corp. Whitepaper： NVIDIA's next generation CUDA compute architecture [R/OL]. [2012?05?18]. http：//www. insidehpc.com.

[2] NVIDIA Corp. Tuning CUDA applications for Fermi [R/OL]. [2011?05?03]. http：// www. people.maths.ox.ac.uk/gilesm/cuda/doc/Fermi_Tuning_Guide.

[3] NVIDIA Corp. CUDA compute unified device architecture， programming guide， Version 3.0 [R/OL]. [2010?05?03]. http：// www. mohamedfahmed.wordpress.com

[4] CUI Xiang， CHEN Yi?feng， MEI Hong， et al. Auto?tuning GEMM for GPGPU with Cache [C]// Proceedings of 2010 IEEE 16th International Conference on Parallel and Distributed Systems （ICPADS）. Shanghai， China： IEEE， 2010： 237?242.

[5] VOLKOV V， DEMMEL J W. Benchmarking GPUs to tune dense linear algebra [C]// Proceedings of 2008. International Conference for High Performance Computing， Networking， Storage and Analysis Austin， TX： [s.n.]， 2008： 1?11.

[6] 李曉雯，崔翔.GPU矩陣乘法和FFT算法的性能優(yōu)化[J].現代電子技術，2013，36（4）：80?84.

1.3 32/64位設備代碼

在Fermi體系結構上，如果代碼按照64?bit的模式編譯，則CUDA編譯器會將CPU代碼和設備代碼都編譯成為64?bit的目標代碼。在這種情況下，設備代碼中的指針變量將會占用多出一倍的寄存器空間。由于Tesla C2050的設備存儲器容量不超過4 GB（在增加ECC的情況下可用的設備存儲器空間只有2.625 GB），因此完全沒有必要在設備代碼中使用64 b的指針。因此，在本文實現的Fermi體系結構上的GEMM代碼中，CPU代碼和設備代碼總是被分別編譯的。

1.4 設備存儲器訪問

在GT200體系結構上，對設備存儲器的訪問是按照半個warp的單位來進行處理的，而在Fermi體系結構上，是按照一個warp的單位來進行處理的。因此，程序員需要調整內核調用時的維度設置。對于具有兩個維度的線程塊，其x維度的大小應該是warp大小的整倍數，而非半個warp大小的整倍數，從而使得每一個warp在訪問設備存儲器時可以得到較高的性能。

1.5 Bank沖突

在GT200體系結構上，共享存儲器具有16個bank，而且對其的訪問是按照半個warp的單位來進行處理的，而在Fermi體系結構上，共享存儲器具有32個bank，而且對其的訪問是按照整個warp的單位來進行處理的。每一個bank的大小為32位。

1.6 多內核并行執(zhí)行

Fermi體系結構支持多個內核代碼的并行執(zhí)行，使得不同應用上下文的內核代碼可以同時在一個GPU上運行；這樣，多個小的內核代碼可以共同利用一個GPU上的計算資源。這也是Fermi體系結構的新特性，但在本文自動優(yōu)化的GEMM代碼中并未使用到。

2 自動優(yōu)化的GEMM代碼

2.1 GEMM代碼模板

文獻[4]描述了在GTX280實現的達到393 Gflops性能的SGEMM內核代碼，在此，依然使用該代碼作為實現自動優(yōu)化的代碼模板。在該實現中，一個Csub被一個線程塊進行計算，根據線程塊中線程數目的多少，一個線程可以計算Csub的半列或多列元素。例如，對于m= 16和n=64，而線程塊具有16×4個線程，則每一個線程將計算一整列Csub的元素；如果線程塊具有16×8個線程，則每一個線程將計算半列Csub的元素。與文獻[4]中的實現類似，每一個線程在使用線程ID計算出訪問矩陣A、B和C的指針位置之后，進入一個循環(huán)。在每一輪循環(huán)中，一個線程塊從設備存儲器中讀入一個Asub的數據到共享存儲器中，之后，又一個內層的循環(huán)被執(zhí)行：在每一輪內層循環(huán)中，一個線程從矩陣B中讀入一個或多個元素，把這些元素與共享存儲器中相應的數據做計算，將結果累加到Csub對應的寄存器中。最后，每個線程將其計算的Csub的數據寫回到設備存儲器中。

盡管與文獻[4]中的GEMM實現具有類似的代碼結構，但是為了滿足緩存的友好性，對該代碼模板進行了一個重要的修正。在此計算中，矩陣B總是位于設備存儲器中，一列線程塊總是需要讀取一列Bsub的數據，如圖2（a）所示。如果線程塊能夠按照列優(yōu)先的順序調度，則可以達到更好的緩存命中效果。因此，在此GEMM模板中，blockIdx.x和blockIdx.y變量的順序被對調，從而達到如圖2（b）所示的設備存儲器訪問效果。這樣，在進行這個轉置之后，線程塊被以較優(yōu)的緩存命中效果的方式調度。在后面給出的性能測試中，將對比這個對調進行或不進行的性能結果。

2.2 對代碼模板進行自動優(yōu)化

設計的自動調節(jié)程序根據代碼模板生成代碼并測試其性能結果。在代碼模板中，5個參數（m，k，n，tx和ty）決定代碼模板的執(zhí)行行為；此外，引入2個額外的參數，一個用來決定線程塊的維度是否對調，來測試緩存的效果，另一個確定L1緩存和共享存儲器的配置比例。因此，整個代碼模板的行為是由7個參數確定的。

圖3顯示經過優(yōu)化得到的矩陣大小為2 048時的DGEMM代碼的實例。注意：在該代碼中，blockIdx.x和blockIdx.y變量通過C語言的宏進行了對調，而相應的調用代碼也做了相應的改變。在此代碼中，m=8，k= 64，n=1 024，tx=64，ty=8，因此一個線程負責計算Csub中的兩列數據。

3 性能測試

4 結 語

為了在Fermi體系結構上書寫高效的代碼，程序員需要很好地了解Fermi體系結構的新特性，以及它們是如何影響程序的性能的。對于Fermi體系結構，程序員尤其要關心緩存對性能的影響。對于Fermi這種復雜的行為和性能難以預測的硬件體系結構，自動優(yōu)化技術不失為得到高性能代碼的一種實用技術。

參考文獻

[1] NVIDIA Corp. Whitepaper： NVIDIA's next generation CUDA compute architecture [R/OL]. [2012?05?18]. http：//www. insidehpc.com.

[2] NVIDIA Corp. Tuning CUDA applications for Fermi [R/OL]. [2011?05?03]. http：// www. people.maths.ox.ac.uk/gilesm/cuda/doc/Fermi_Tuning_Guide.

[3] NVIDIA Corp. CUDA compute unified device architecture， programming guide， Version 3.0 [R/OL]. [2010?05?03]. http：// www. mohamedfahmed.wordpress.com

[4] CUI Xiang， CHEN Yi?feng， MEI Hong， et al. Auto?tuning GEMM for GPGPU with Cache [C]// Proceedings of 2010 IEEE 16th International Conference on Parallel and Distributed Systems （ICPADS）. Shanghai， China： IEEE， 2010： 237?242.

[5] VOLKOV V， DEMMEL J W. Benchmarking GPUs to tune dense linear algebra [C]// Proceedings of 2008. International Conference for High Performance Computing， Networking， Storage and Analysis Austin， TX： [s.n.]， 2008： 1?11.

[6] 李曉雯，崔翔.GPU矩陣乘法和FFT算法的性能優(yōu)化[J].現代電子技術，2013，36（4）：80?84.