OpenCL計算軟件棧評估*

朱 浩，周博洋，盧雪山，杜溢墨

(1.軍事科學(xué)院國防科技創(chuàng)新研究院，北京 100000；2.國防科技大學(xué)計算機(jī)學(xué)院，湖南 長沙 410073；3.空軍后勤部，北京100000；4.31008部隊，北京100091)

1 引言

人工智能及其支撐技術(shù)將成為決定未來社會發(fā)展的重要因素之一，人工智能技術(shù)在感知與信息處理、指揮決策、網(wǎng)絡(luò)空間安全和無人系統(tǒng)等領(lǐng)域正發(fā)揮著越來越重要的作用。以使用GPU為代表的異構(gòu)計算進(jìn)行人工智能應(yīng)用處理的效能更高，因此研究人員提出了CPU+GPU的異構(gòu)模式，由CPU進(jìn)行邏輯控制和數(shù)據(jù)轉(zhuǎn)發(fā)，而GPU則負(fù)責(zé)大規(guī)模運(yùn)算，這樣的各取所長的工作方式更有利于GPU性能的發(fā)揮。

計算軟件棧是發(fā)揮GPU硬件性能的關(guān)鍵。NVIDIA公司專門針對自己生產(chǎn)的系列GPU設(shè)計了CUDA計算軟件棧。CUDA提供了簡單的開發(fā)工具可用于設(shè)計GPU應(yīng)用程序，降低了開發(fā)難度，并與計算庫、驅(qū)動等相配合實現(xiàn)了程序在GPU上高效的在線編譯、提交和運(yùn)行。與CUDA相比，開放運(yùn)算語言O(shè)penCL(Open Computing Language)[1]是一個面向異構(gòu)系統(tǒng)的完全免費(fèi)的通用標(biāo)準(zhǔn)，并且適用于多種架構(gòu)的處理器。OpenCL僅僅是一個通用的API，只提供函數(shù)接口，在此基礎(chǔ)上衍生出了很多開源的計算軟件棧，比如AMD維護(hù)的ROCm[2]、Freedesktop維護(hù)的Mesa[3]等，便于人們應(yīng)用和研究。

雖然CUDA擁有比較好的性能和市場表現(xiàn)，但是由于目前CUDA源碼閉源并且主要支持NVIDIA的GPU和Intel商用平臺，沒有廣泛的移植性?？紤]到未來在飛騰CPU和麒麟操作系統(tǒng)等國產(chǎn)軟硬件平臺上會有大量GPU計算移植和適配的需求[4,5]，本文重點(diǎn)研究開源的OpenCL計算軟件棧包括Mesa、ROCm(Radeon Open Computing)等，研究評估不同計算軟件棧在程序各個執(zhí)行階段的性能差異、在Intel平臺和國產(chǎn)平臺上的性能差異等，為OpenCL軟件棧的選型和優(yōu)化提供決策依據(jù)。

2 OpenCL計算軟件棧

2.1 OpenCL

OpenCL是第一個面向CPUGPUFPGA等異構(gòu)系統(tǒng)并行編程的開放式、免費(fèi)標(biāo)準(zhǔn)，也是一個統(tǒng)一的編程環(huán)境，當(dāng)前最新協(xié)議為OpenCL 3.0。OpenCL協(xié)議有多種系統(tǒng)實現(xiàn)，目前主流計算軟件棧有Mesa和ROCm。

OpenCL應(yīng)用的流程包含編譯、數(shù)據(jù)交換和執(zhí)行等多個過程，如圖1所示，具體到源碼通常包含clGetPlatformIDs(獲取平臺)、clGetDeviceIDs(獲取設(shè)備)、clCreateContext(創(chuàng)建Context環(huán)境)、clCreateCommandQueue(創(chuàng)建命令隊列)、clCreateBuffer(創(chuàng)建緩沖對象)、clEnqueueWriteBuffer(Host寫數(shù)據(jù)到GPU)、clCreateProgramWithSource(創(chuàng)建程序，為Context環(huán)境創(chuàng)建程序?qū)ο?，并將Kernel源碼加載到該對象中)、clBuildProgram(編譯程序)、clCreateKernel(創(chuàng)建Kernel)、clSetKernelArg(設(shè)置Kernel參數(shù))、clEnqueueNDRangeKernel(提交并執(zhí)行Kernel)和clEnqueueReadBuffer(數(shù)據(jù)拷貝回Host)等函數(shù)，OpenCL提供了這些函數(shù)接口的定義。而Mesa和ROCm等軟件棧的主要區(qū)別就在于以不同的方式實現(xiàn)了這些函數(shù)，所以性能會有差異。

Figure 1 Running process of OpenCL applications圖1 OpenCL應(yīng)用的運(yùn)行流程

OpenCL應(yīng)用程序包含C程序代碼和Kernel程序代碼。C程序代碼負(fù)責(zé)除內(nèi)核計算以外的設(shè)備初始化、數(shù)據(jù)輸入輸出、Kernel程序提交和執(zhí)行等部分；Kernel程序代碼是在GPU上執(zhí)行的計算部分。OpenCL應(yīng)用程序的內(nèi)核部分是通過LLVM(Low Level Virtual Machine)編譯器[6]運(yùn)行時編譯(Just In Time Compile)，LLVM編譯器前端負(fù)責(zé)對程序進(jìn)行語法分析、詞法分析以及優(yōu)化后轉(zhuǎn)成LLVM IR格式的中間語言；LLVM后端根據(jù)當(dāng)前使用何種GPU設(shè)備來完成中間語言到GPU可執(zhí)行語言的編譯。

2.2 Mesa

Mesa(全稱Mesa 3D)是一個完全開源的設(shè)備驅(qū)動程序，支持多種GPU，于1993年8月由布萊恩·保羅提出，設(shè)計之初僅僅是為了實現(xiàn)OpenGL的功能。在隨后的發(fā)展中，其內(nèi)容不斷豐富并加入了GLX(OpenGL extension to X)、EGL、OpenCL等功能的實現(xiàn)。其中實現(xiàn)OpenCL語言的項目叫做Clover。

Mesa Clover能夠支持OpenCL 2.2協(xié)議，對AMD 顯卡有良好的支持。Mesa Clover結(jié)構(gòu)簡單，在處理數(shù)據(jù)量不大的OpenCL應(yīng)用時，相比其它軟件棧有一定的性能優(yōu)勢。

2.3 ROCm

ROCm是AMD維護(hù)的成熟的完全開源的OpenCL實現(xiàn)，支持OpenCL 2.2。ROCm有比較完整的生態(tài)，不僅支持OpenCL計算，還可以通過HIP(Heterogenous-compute Interface for Portability)支持CUDA計算，還包含豐富的調(diào)試、性能分析和開發(fā)環(huán)境等工具，是GPU計算方向比較活躍的項目。

ROCm主要包括3個部分：ROCm 庫、ROCm平臺和ROCm內(nèi)核驅(qū)動，如圖2所示。

Figure 2 Architecture of ROCm圖2 ROCm的軟件棧結(jié)構(gòu)

ROCm庫包含OpenCL API和OpenCL庫，其中OpenCL API是規(guī)范中定義的標(biāo)準(zhǔn)接口，ROCm OpenCL庫就是OpenCL API的實現(xiàn)。ROCm平臺主要實現(xiàn)了異構(gòu)系統(tǒng)的功能，從系統(tǒng)層面整合CPU和GPU資源，也支持AMD的獨(dú)立GPU計算。最大特點(diǎn)是對異構(gòu)資源統(tǒng)一定義了虛擬地址空間，通過共享指針來共享數(shù)據(jù)，減少數(shù)據(jù)在不同資源間的拷貝次數(shù)。ROCm內(nèi)核驅(qū)動用來支持異構(gòu)系統(tǒng)的內(nèi)核驅(qū)動，其主要功能包括支持任務(wù)隊列，簡化計算任務(wù)在CPU和GPU上的分布，支持異構(gòu)存儲管理以及驅(qū)動間的交互。其中，AMDKFD主要負(fù)責(zé)異構(gòu)計算任務(wù)的處理，AMDGPU Dirver除了支持顯示和計算以外，還提供GPU設(shè)備的初始化，為AMDKFD提供設(shè)備信息等功能。

3 國產(chǎn)平臺和商用平臺對比評估

本節(jié)對國產(chǎn)飛騰平臺和Intel商用平臺上的軟件棧性能進(jìn)行對比分析，了解CPU性能對GPU應(yīng)用計算的影響程度。進(jìn)行平臺對比測試的環(huán)境如表1所示，本節(jié)采用了飛騰1500A 4核處理器(后文稱FT 1500A)和Intel i7 4核處理器，選擇了當(dāng)前飛騰臺式機(jī)上主流、適配成熟的顯卡執(zhí)行計算，OpenCL軟件棧為Mesa 17.0。

Table 1 Test environment of Pytium and Intel platforms表1 國產(chǎn)飛騰平臺和商用Intel平臺軟件棧性能測試環(huán)境

OpenCL測試程序采用Black-Scholes基準(zhǔn)程序(簡稱B-S)，該程序用于計算期權(quán)定價，數(shù)據(jù)輸入是期權(quán)數(shù)集合(為了對比更加直觀，本文將存儲浮點(diǎn)型期權(quán)占用的存儲空間轉(zhuǎn)化成數(shù)據(jù)量的大小)，直接反映的是數(shù)據(jù)并行處理能力。對Black-Scholes基準(zhǔn)程序的編譯、數(shù)據(jù)讀寫、運(yùn)行和完成等過程進(jìn)行采樣，測試這些OpenCL接口的執(zhí)行時間以及GPU加速比(CPU計算時間/GPU計算時間)。后面各圖中X軸是將期權(quán)數(shù)轉(zhuǎn)化成數(shù)據(jù)量大小后的值，Y軸是執(zhí)行時間。

如圖3所示，Intel i7處理器的應(yīng)用程序內(nèi)核編譯性能是FT 1500A處理器的5倍左右，由于是采用多線程編譯，這個數(shù)據(jù)可以反映出不修改LLVM編譯器默認(rèn)條件下Intel i7處理器的多核并行計算能力大概是FT 1500A的5倍。

Figure 3 Build kernel time of B-S application on Phytium and Intel platforms圖3 B-S應(yīng)用在飛騰和Intel平臺上的編譯內(nèi)核時間

Figure 4 Read buffer time of B-S application on Phytium and Intel platforms圖4 B-S應(yīng)用在飛騰和Intel平臺上的Read buffer時間

Read buffer時間反映的是CPU從GPU讀取處理后數(shù)據(jù)的能力。如圖4所示，當(dāng)數(shù)據(jù)量小于175 MB時，F(xiàn)T 1500A執(zhí)行時間大約是Intel i7的3.5倍；當(dāng)數(shù)據(jù)量大于175 MB時，F(xiàn)T 1500A執(zhí)行時間大約是Intel i7的2倍。

評估內(nèi)核函數(shù)在GPU上的實際執(zhí)行時間，記錄為GPU計算時間。如圖5所示，開始FT 1500A上GPU計算時間幾乎與Intel i7的相等，數(shù)據(jù)量在240 MB附近時，F(xiàn)T 1500A的GPU計算時間出現(xiàn)跳躍式大幅度增加，接近4倍，而數(shù)據(jù)量在240 MB附近時，Intel i7的GPU計算時間有小幅度增加。考慮到這部分時間主要運(yùn)行在GPU端，導(dǎo)致時間的差異應(yīng)該是與CPU性能無關(guān)的，經(jīng)過分析發(fā)現(xiàn)主要是由于CPU體系結(jié)構(gòu)的不同，導(dǎo)致編譯后的內(nèi)核代碼不同。在對內(nèi)核代碼進(jìn)行編譯時，LLVM編譯器會根據(jù)CPU體系結(jié)構(gòu)的不同，編譯生成字的節(jié)碼有所區(qū)別。在Intel商用平臺上使用其他處理器運(yùn)行B-S應(yīng)用來測試GPU計算時間，確實沒有區(qū)別，這證實了主要跟體系結(jié)構(gòu)有關(guān)。

Figure 5 GPU running time of B-S application on Phytium and Intel platforms圖5 B-S應(yīng)用在飛騰和Intel平臺上GPU計算時間

為了了解是國產(chǎn)飛騰平臺還是商用Intel平臺上加速效果更明顯，獲得GPU執(zhí)行計算的加速比，首先將B-S應(yīng)用內(nèi)核函數(shù)的算法重寫并使用CPU進(jìn)行計算，得到2個平臺上各自的CPU計算時間。如圖6a所示，Intel i7的性能優(yōu)勢隨著數(shù)據(jù)量的增大越來越明顯，F(xiàn)T 1500A的執(zhí)行時間最大時大概是Intel i7的2.8倍。用CPU計算時間除以GPU計算時間和可以得到應(yīng)用程序的加速比。當(dāng)數(shù)據(jù)量低于240 MB左右時，F(xiàn)T 1500A上的加速比高于Intel i7上的，超過20倍。當(dāng)數(shù)據(jù)量超過240 MB時，Intel i7上的加速比高于FT 1500A上的，這是因為在FT 1500A上，內(nèi)核的GPU執(zhí)行時間在數(shù)據(jù)量超過240 MB時出現(xiàn)跳躍點(diǎn)，性能突然變差(如圖6b所示)。整體來說，Intel i7上的加速比大概在5～10倍波動，F(xiàn)T 1500A上的加速比大概在4～20倍波動，最后兩者分別穩(wěn)定在5倍左右。

Figure 6 CPU running time of B-S application on Phytium and Intel platforms圖6 B-S應(yīng)用在飛騰和Intel平臺上的CPU計算時間

Figure 7 Overall running time of B-S application on Phytium and Intel platforms圖7 B-S應(yīng)用在飛騰和Intel平臺上的運(yùn)行時間

如圖7a所示，考慮測試初始化、編譯、數(shù)據(jù)準(zhǔn)備和GPU計算等過程的運(yùn)行總時間，F(xiàn)T 1500A上的執(zhí)行時間是Intel i7上的3.3～5.2倍。如果不考慮初次運(yùn)行的編譯時間，如圖7b所示，F(xiàn)T 1500A上的執(zhí)行時間是Intel i7上的2～4.1倍。

結(jié)合以上測試數(shù)據(jù)，對在國產(chǎn)平臺和商用平臺上進(jìn)行GPU計算的性能評估總結(jié)如下：

(1)OpenCL應(yīng)用的編譯時間、數(shù)據(jù)拷貝時間、GPU計算時間等，也就是GPU應(yīng)用的整體性能都和CPU平臺的性能、體系結(jié)構(gòu)有關(guān)。國產(chǎn)平臺和商用平臺2種不同CPU體系結(jié)構(gòu)上的GPU計算時間(即內(nèi)核函數(shù)在GPU上的執(zhí)行時間)會有區(qū)別，主要是因為內(nèi)核程序編譯生成的可執(zhí)行代碼受體系結(jié)構(gòu)的影響有差異。

(2)在計算數(shù)據(jù)量較小時，F(xiàn)T 1500A上使用GPU計算獲得的加速比更大；數(shù)據(jù)量越大尤其是超過某個值時，GPU計算時間會有不同程度的跳漲，Intel平臺上的GPU加速性能更優(yōu)。這個跳漲的時機(jī)和原因有待進(jìn)一步分析。當(dāng)要處理的數(shù)據(jù)量較大時，國產(chǎn)平臺上受制于CPU性能和體系結(jié)構(gòu)編譯優(yōu)化方式等瓶頸因素，GPU計算獲取的加速收益不如Intel商用平臺。

(3)雖然大數(shù)據(jù)量計算時國產(chǎn)平臺上計算加速收益不如Intel商用平臺，但如果不考慮編譯時間，F(xiàn)T 1500A上OpenCL應(yīng)用執(zhí)行時間比Intel i7 CPU上長2～4.1倍，小于FT 1500A與Intel i7 CPU本身的性能差距(這個差距大概是5倍)。在有的程序既可以使用CPU計算也可以使用GPU計算時，在國產(chǎn)平臺上應(yīng)更多地使用GPU進(jìn)行計算，減少CPU性能差異帶來的負(fù)面影響。

(4)在國產(chǎn)平臺上優(yōu)化OpenCL計算，應(yīng)該是整個軟硬件系統(tǒng)的優(yōu)化，而不能僅僅優(yōu)化GPU硬件體系結(jié)構(gòu)。

4 軟件棧對比評估

Mesa和ROCm是當(dāng)下2個比較熱門的OpenCL開源計算軟件棧，本節(jié)對比了OpenCL應(yīng)用在2個軟件棧上的性能表現(xiàn)，作為對照還增加了AMD-APPSDK軟件棧的測試，APPSDK是AMD公司維護(hù)的一個閉源的OpenCL計算軟件棧?？紤]不同平臺上的驅(qū)動和內(nèi)核有差異，可能會對應(yīng)用執(zhí)行產(chǎn)生影響，最后對驅(qū)動、內(nèi)核造成的運(yùn)行時間影響進(jìn)行評估。

4.1 OpenCL軟件棧測試

OpenCL軟件棧測試環(huán)境如表2所示，采用了Intel i7處理器平臺，裝配有AMD中端檔次的RX460顯卡，內(nèi)核版本是4.15，驅(qū)動分別是開源的amdgpu和閉源的amdgpu pro，執(zhí)行B-S應(yīng)用。測試時對比的OpenCL軟件棧分別為ROCm、APPSDK、Mesa-open(Mesa，采用開源驅(qū)動)和Mesa-close(Mesa，采用閉源驅(qū)動)。

Table 2 Test environment of software stack 表2 軟件棧測試環(huán)境

如圖8所示，進(jìn)行Create Buffer緩存的分配時，ROCm的執(zhí)行時間隨著數(shù)據(jù)量的增長而快速上升，其他3種軟件棧幾乎沒有變化，多次測試均是這個結(jié)果。通過對代碼分析發(fā)現(xiàn)，ROCm在緩存分配時會進(jìn)行實際的物理分配，而其他的軟件棧只是創(chuàng)建一個mem結(jié)構(gòu)體，在讀寫時才進(jìn)行實際的物理空間分配。

Figure 8 Create buffer running time of B-S application on various software stacks圖8 B-S應(yīng)用在不同軟件棧上的Create Buffer執(zhí)行時間

如圖9所示，對比內(nèi)核程序編譯的時間，Mesa上編譯時間最長，其次是ROCm，最短的是APPSDK。經(jīng)過分析發(fā)現(xiàn)，編譯性能的差異主要是由于LLVM編譯器在對內(nèi)核源碼進(jìn)行編譯生成ISA代碼過程中，編譯后端的目標(biāo)組合(Target Triple)差異性導(dǎo)致的，比如Mesa的目標(biāo)組合是(amdgcn，Mesa3d，Mesa3d)，ROCm的目標(biāo)組合是(amdgcn，amd，amdhsa)。這個編譯目標(biāo)組合是在LLVM編譯器里定義的，整個三元組含義是：(AMD GPU體系結(jié)構(gòu)，廠商，操作系統(tǒng)環(huán)境)。根據(jù)目標(biāo)組合的不同，不同的軟件棧編譯成的ISA代碼就會有差異，從而導(dǎo)致編譯內(nèi)核執(zhí)行時間有所差距，最終導(dǎo)致在GPU運(yùn)行內(nèi)核程序的效率有差別。

Figure 9 Build kernel running time of B-S application on various software stacks圖9 B-S應(yīng)用在不同軟件棧上的編譯內(nèi)核執(zhí)行時間

如圖10所示，內(nèi)核程序在GPU上計算時，在數(shù)據(jù)量低于240 MB時,APPSDK和ROCm的執(zhí)行時間都略高于Mesa的。對Mesa來說，GPU上的計算時間仍存在一個跳躍點(diǎn)，當(dāng)數(shù)據(jù)量高于240 MB時，Mesa開閉源的性能都會突然變差，執(zhí)行時間達(dá)到APPSDK和ROCm的2倍。

Figure 10 GPU running time of B-S application on various software stacks圖10 B-S應(yīng)用在不同軟件棧上的GPU計算時間

如圖11a所示，將主機(jī)內(nèi)存數(shù)據(jù)寫入GPU存儲時,Mesa軟件棧無論開閉源驅(qū)動，它們的執(zhí)行性能都相近，都低于ROCm和APPSDK上的性能。但是，當(dāng)數(shù)據(jù)量高于240 MB時，存在一個跳躍點(diǎn)，Mesa開源驅(qū)動上的執(zhí)行時間長于Mesa閉源驅(qū)動的。

如圖11b所示，將GPU計算產(chǎn)生的數(shù)據(jù)讀回主機(jī)內(nèi)存上是Read buffer，與Write buffer的表現(xiàn)類似,Mesa的性能差于APPSDK和ROCm的。在數(shù)據(jù)量為240 MB左右時，Mesa的執(zhí)行時間有個快速增長的區(qū)間。

Figure 11 Read/Write buffer time of B-S application on various software stacks圖11 B-S應(yīng)用在不同軟件棧上的讀寫緩沖執(zhí)行時間

如圖12所示，對比B-S應(yīng)用從提交到執(zhí)行結(jié)束的總執(zhí)行時間，APPSDK有最好的性能表現(xiàn)，當(dāng)數(shù)據(jù)量較小時(100 MB左右)，ROCm上的執(zhí)行時間長于Mesa上的。但是，當(dāng)數(shù)據(jù)量在240 MB附近時,Mesa上的執(zhí)行時間會快速增加。

Figure 12 Overall running time of B-S application on various software stacks圖12 B-S應(yīng)用在不同軟件棧上的總時間

對OpenCL軟件棧測試總結(jié)如下：

(1)應(yīng)用總的執(zhí)行時間，最優(yōu)的是APPSDK，ROCm次之，然后是Mesa加閉源驅(qū)動的組合，最后是Mesa加開源驅(qū)動的組合。

(2)對GPU計算時間來說，尤其是大數(shù)據(jù)量時，APPSDK、ROCm平臺的GPU性能優(yōu)于Mesa的。APPSDK和ROCm的性能曲線是基本相同的，Mesa平臺開閉源性能曲線基本相同。這說明了LLVM編譯器針對APPSDK和ROCm的優(yōu)化更好，相同內(nèi)核函數(shù)在這兩者上的執(zhí)行效率更高。

從圖10可以看出GPU計算時間，Mesa性能明顯差于APPSDK和ROCm的，并且在數(shù)據(jù)量為240 MB時會有一個跳躍點(diǎn)，執(zhí)行時間會明顯增長。為了證明普遍性，本文又對比了另外幾個應(yīng)用的GPU計算性能。

如圖13所示，進(jìn)行2D圖像卷積(Convolution-2D)測試，在Mesa上的GPU計算時間上升幅度明顯高于ROCm上的，當(dāng)數(shù)據(jù)量達(dá)到240 MB附近時，也出現(xiàn)跳躍式增加。而ROCm上的GPU計算時間一直平穩(wěn)上升，在數(shù)據(jù)量達(dá)到280 MB時，增速僅僅有小幅度增加。這說明對于Mesa，數(shù)據(jù)量達(dá)到一定閾值時，應(yīng)用的性能會有較大衰減，原因初步分析與Mesa中固定內(nèi)存的創(chuàng)建方式有關(guān)，更深入的研究是下一步的工作。

Figure 13 GPU running time of 2D convolution application on various ROCm and Mesa-open圖13 ROCm和Mesa-open上2D圖像卷積測試的GPU計算時間

Rodinia srad是超聲波和雷達(dá)成像應(yīng)用，用于消除局部相關(guān)噪聲，而不會破壞重要的圖像特征。Rodinia hotspot用于根據(jù)建筑平面布置圖和模擬功率測量來估算處理器溫度。如圖14所示，Mesa上的Rodinia srad GPU計算時間為ROCm上的1.97倍。對于Rodinia hotspot應(yīng)用，Mesa的GPU計算時間為ROCm的3倍。進(jìn)一步證實了Mesa平臺上的GPU計算性能確實不如ROCm上的。

Figure 14 GPU running time of Rodinia applications on ROCm and Mesa-open圖14 Rodina應(yīng)用在ROCm和Mesa-open棧上的GPU計算時間

下面2小節(jié)對驅(qū)動、內(nèi)核差異造成的運(yùn)行時間影響進(jìn)行評估。OpenCL應(yīng)用運(yùn)行依賴的OpenCL軟件棧還包括系統(tǒng)軟件棧，比如底層的驅(qū)動和內(nèi)核，它們構(gòu)成了一個完整的軟件棧來實現(xiàn)應(yīng)用與硬件的交互。

4.2 開閉源驅(qū)動測試

從圖12可以看出，應(yīng)用在Mesa加開源驅(qū)動上運(yùn)行與Mesa加閉源驅(qū)動的相比，性能要差一些，并且隨著數(shù)據(jù)量的增大，性能差距越來越大。為了說明這種現(xiàn)象的普遍性，本節(jié)選擇使用另一種OpenCL應(yīng)用Adi進(jìn)行確認(rèn)測試。Adi是針對多核CPU和GPU的基準(zhǔn)測試套件PolyBench-ACC中的測試用例。

如圖15所示，與B-S應(yīng)用的Write buffer和Read buffer類似，Adi應(yīng)用在Mesa開源驅(qū)動上運(yùn)行與Mesa閉源驅(qū)動上相比，開源驅(qū)動的性能要差，當(dāng)數(shù)據(jù)量小時差距比較小，開源驅(qū)動上的運(yùn)行時間波動大。

Figure 15 Read/Write buffer time of Adi application on various software stacks圖15 Adi應(yīng)用在不同軟件棧上的讀寫緩沖執(zhí)行時間

如圖16所示，與B-S應(yīng)用類似，Mesa上的Adi應(yīng)用GPU計算性能比較差，并且開閉源驅(qū)動均會在240 MB數(shù)據(jù)量時出現(xiàn)較大波動。

Figure 16 GPU running time of Adi application圖16 Adi應(yīng)用在GPU上的計算時間

根據(jù)B-S應(yīng)用和Adi應(yīng)用等的性能表現(xiàn)，總結(jié)開閉源軟件棧測試結(jié)果如下：

(1)開閉源驅(qū)動會對讀寫緩沖和GPU計算時間等過程產(chǎn)生影響。

(2)在數(shù)據(jù)量未超過跳躍點(diǎn)時，開閉源驅(qū)動執(zhí)行的總時間幾乎相等，開閉源驅(qū)動的GPU性能沒有差異，當(dāng)數(shù)據(jù)量超過跳躍點(diǎn)后，閉源驅(qū)動的GPU性能略高于開源驅(qū)動。經(jīng)過對同一時間點(diǎn)開發(fā)下載的開閉源驅(qū)動對比分析發(fā)現(xiàn)，無論是OpenCL軟件棧還是驅(qū)動閉源軟件棧的版本，相對于開源版本都比較新，廠商會優(yōu)先將一些新的特性優(yōu)化集成在閉源版本中，這是GPU廠商的市場策略，這也是閉源性能有優(yōu)勢的主要原因。

4.3 內(nèi)核版本測試

這一節(jié)評估內(nèi)核版本對B-S應(yīng)用進(jìn)行GPU加速計算的影響。內(nèi)核V4.15是Ubuntu1804版本默認(rèn)采用的內(nèi)核，與之前的長期支持LTS內(nèi)核版本有較大的改動，本節(jié)測試選擇了在4.15前后的2個內(nèi)核進(jìn)行對比，即Ubuntu1604.2默認(rèn)的4.8內(nèi)核版本與Ubuntu1804.1默認(rèn)的 4.18版本。

如圖17所示，在數(shù)據(jù)量小于100 MB時，高低版本總執(zhí)行時間接近，隨著數(shù)據(jù)量增大，4.8版本的總執(zhí)行時間相比4.18版本的總執(zhí)行時間越來越長，超過300 MB時，4.8版本的總執(zhí)行時間接近4.18版本的2倍。

Figure 17 Overall running time of B-S application on various kernels圖17 B-S應(yīng)用在不同內(nèi)核條件下的總運(yùn)行時間

總結(jié)對內(nèi)核版本差異的測試，可以看出：內(nèi)核版本變化對GPU計算時間影響比較微小，主要是在數(shù)據(jù)量達(dá)到一定閾值后，2個版本上的時間才有差別，新版本上的性能較優(yōu)。

4.4 總結(jié)

對第4節(jié)的測試結(jié)果總結(jié)如下:

(1)Mesa平臺上的GPU應(yīng)用有一個跳躍點(diǎn)(數(shù)據(jù)量240 MB左右)，超過這個點(diǎn)，應(yīng)用執(zhí)行時間會跳躍式增長；但是ROCm不存在這種明顯的跳躍點(diǎn)，執(zhí)行時間增長比較平穩(wěn)。Mesa與ROCm的GPU計算性能差異與算法和數(shù)據(jù)量有關(guān)，ROCm的GPU計算性能為Mesa的2.5～15倍左右(算法越復(fù)雜，數(shù)據(jù)量越大，ROCm優(yōu)勢越大)。

(2)ROCm開源平臺相對于APPSDK閉源平臺，應(yīng)用總的執(zhí)行時間在數(shù)據(jù)量小時存在劣勢，存在很大的優(yōu)化空間。相比Mesa，ROCm是以后在國產(chǎn)平臺上進(jìn)行移植和適配優(yōu)化比較理想的GPU計算軟件棧。

(3)在閉源驅(qū)動、新版本的驅(qū)動和內(nèi)核上運(yùn)行GPU應(yīng)用運(yùn)行性能有一定的優(yōu)勢，在國產(chǎn)平臺上進(jìn)行GPU計算時盡可能選擇新版本的操作系統(tǒng)。

5 相關(guān)工作

Karimi等人[7-9]評估了NVIDIA GPU上的OpenCL和CUDA性能，對比了數(shù)據(jù)遷移性能、內(nèi)核執(zhí)行時間和總的執(zhí)行時間，CUDA在某些場景下相對OpenCL有比較好的性能，但OpenCL的高可移植性決定了它是CUDA的一個很好替代。Komatsu等人[10,11]評估了OpenCL應(yīng)用的性能和可移植性，OpenCL應(yīng)用可以通過調(diào)優(yōu)達(dá)到理想的性能。Mukherjee等人[12]對比了2種軟件棧OpenCL和異構(gòu)系統(tǒng)體系結(jié)構(gòu)HSA(Heterogeneous System Architecture)，HSA是HSA基金會維護(hù)的一種異構(gòu)計算軟件棧，由于HSA提供了共享虛擬內(nèi)存、動態(tài)并行化等新特性，在異構(gòu)平臺上相比OpenCL具有更好的性能。另外還有一些研究是關(guān)于OpenCL測試基準(zhǔn)的設(shè)計和實現(xiàn)[13-15]，用于分析異構(gòu)系統(tǒng)的瓶頸和優(yōu)化系統(tǒng)性能。本文主要是在國產(chǎn)和商用平臺上對不同OpenCL軟件棧的性能進(jìn)行對比分析，從系統(tǒng)軟件棧和系統(tǒng)優(yōu)化角度進(jìn)行更加全面深入的研究。

6 結(jié)束語

本文基于Black-Scholes、2D-Convolution、Rodinia等幾種測試集評估OpenCL軟件棧在不同條件下的性能表現(xiàn)，為GPU計算機(jī)軟件棧選型和優(yōu)化提供依據(jù)。通過對國產(chǎn)和商用平臺上OpenCL性能測試發(fā)現(xiàn)，CPU對GPU計算有比較大的影響，對GPU計算的優(yōu)化應(yīng)該是從CPU到GPU、從軟件到硬件成體系的優(yōu)化。通過對比和測試不同軟件棧發(fā)現(xiàn)，ROCm的性能和穩(wěn)定性都優(yōu)于Mesa等軟件棧，比較適合以后在國產(chǎn)平臺上進(jìn)行移植和優(yōu)化。閉源版本和新版本的軟件棧相對于穩(wěn)定的開源版本會有一定的性能優(yōu)勢，為版本選型提供了依據(jù)。