李曉雯　崔翔　殷瑞杰　劉強(qiáng)

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

關(guān)鍵詞： GPU程序設(shè)計(jì)； 矩陣乘法； 自動優(yōu)化； GEMM模板

中圖分類號： TN40； TP312 文獻(xiàn)標(biāo)識碼： 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體系結(jié)構(gòu)相比，新型的Tesla 2050 GPU具有一些新的特性：如增強(qiáng)的雙精度浮點(diǎn)性能、L1/L2 緩存結(jié)構(gòu)、更多的寄存器、更大的共享存儲器、ECC支持和更快的原子操作[1?3]。

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

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

1 Fermi的新特性

1.1 L1/L2緩存

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

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

1.2 寄存器文件

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

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

1.3 32/64位設(shè)備代碼

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

1.4 設(shè)備存儲器訪問

在GT200體系結(jié)構(gòu)上，對設(shè)備存儲器的訪問是按照半個(gè)warp的單位來進(jìn)行處理的，而在Fermi體系結(jié)構(gòu)上，是按照一個(gè)warp的單位來進(jìn)行處理的。因此，程序員需要調(diào)整內(nèi)核調(diào)用時(shí)的維度設(shè)置。對于具有兩個(gè)維度的線程塊，其x維度的大小應(yīng)該是warp大小的整倍數(shù)，而非半個(gè)warp大小的整倍數(shù)，從而使得每一個(gè)warp在訪問設(shè)備存儲器時(shí)可以得到較高的性能。

1.5 Bank沖突

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

1.6 多內(nèi)核并行執(zhí)行

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

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

2.1 GEMM代碼模板

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

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

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

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

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

3 性能測試

4 結(jié) 語

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

參考文獻(xiàn)

[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].現(xiàn)代電子技術(shù)，2013，36（4）：80?84.

1.3 32/64位設(shè)備代碼

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

1.4 設(shè)備存儲器訪問

在GT200體系結(jié)構(gòu)上，對設(shè)備存儲器的訪問是按照半個(gè)warp的單位來進(jìn)行處理的，而在Fermi體系結(jié)構(gòu)上，是按照一個(gè)warp的單位來進(jìn)行處理的。因此，程序員需要調(diào)整內(nèi)核調(diào)用時(shí)的維度設(shè)置。對于具有兩個(gè)維度的線程塊，其x維度的大小應(yīng)該是warp大小的整倍數(shù)，而非半個(gè)warp大小的整倍數(shù)，從而使得每一個(gè)warp在訪問設(shè)備存儲器時(shí)可以得到較高的性能。

1.5 Bank沖突

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

1.6 多內(nèi)核并行執(zhí)行

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

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

2.1 GEMM代碼模板

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

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

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

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

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

3 性能測試

4 結(jié) 語

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

參考文獻(xiàn)

[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].現(xiàn)代電子技術(shù)，2013，36（4）：80?84.

1.3 32/64位設(shè)備代碼

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

1.4 設(shè)備存儲器訪問

在GT200體系結(jié)構(gòu)上，對設(shè)備存儲器的訪問是按照半個(gè)warp的單位來進(jìn)行處理的，而在Fermi體系結(jié)構(gòu)上，是按照一個(gè)warp的單位來進(jìn)行處理的。因此，程序員需要調(diào)整內(nèi)核調(diào)用時(shí)的維度設(shè)置。對于具有兩個(gè)維度的線程塊，其x維度的大小應(yīng)該是warp大小的整倍數(shù)，而非半個(gè)warp大小的整倍數(shù)，從而使得每一個(gè)warp在訪問設(shè)備存儲器時(shí)可以得到較高的性能。

1.5 Bank沖突

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

1.6 多內(nèi)核并行執(zhí)行

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

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

2.1 GEMM代碼模板

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

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

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

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

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

3 性能測試

4 結(jié) 語

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

參考文獻(xiàn)

[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].現(xiàn)代電子技術(shù)，2013，36（4）：80?84.