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

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

1 引言

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

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

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

2 OpenCL計(jì)算軟件棧

2.1 OpenCL

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

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

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

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

2.2 Mesa

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

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

2.3 ROCm

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

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

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

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

3 國(guó)產(chǎn)平臺(tái)和商用平臺(tái)對(duì)比評(píng)估

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

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

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

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

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

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

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

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

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

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

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

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

結(jié)合以上測(cè)試數(shù)據(jù)，對(duì)在國(guó)產(chǎn)平臺(tái)和商用平臺(tái)上進(jìn)行GPU計(jì)算的性能評(píng)估總結(jié)如下：

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

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

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

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

4 軟件棧對(duì)比評(píng)估

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

4.1 OpenCL軟件棧測(cè)試

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

如圖13所示，進(jìn)行2D圖像卷積(Convolution-2D)測(cè)試，在Mesa上的GPU計(jì)算時(shí)間上升幅度明顯高于ROCm上的，當(dāng)數(shù)據(jù)量達(dá)到240 MB附近時(shí)，也出現(xiàn)跳躍式增加。而ROCm上的GPU計(jì)算時(shí)間一直平穩(wěn)上升，在數(shù)據(jù)量達(dá)到280 MB時(shí)，增速僅僅有小幅度增加。這說(shuō)明對(duì)于Mesa，數(shù)據(jù)量達(dá)到一定閾值時(shí)，應(yīng)用的性能會(huì)有較大衰減，原因初步分析與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圖像卷積測(cè)試的GPU計(jì)算時(shí)間

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

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

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

4.2 開閉源驅(qū)動(dòng)測(cè)試

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

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

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

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

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

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

(1)開閉源驅(qū)動(dòng)會(huì)對(duì)讀寫緩沖和GPU計(jì)算時(shí)間等過(guò)程產(chǎn)生影響。

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

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

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

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

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

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

4.4 總結(jié)

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

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

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

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

5 相關(guān)工作

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

6 結(jié)束語(yǔ)

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