包達(dá)爾罕,高文煒,鄭 欣,馮 路,楊金穎
(1.西安微電子技術(shù)研究所,陜西 西安 710054;2.火箭軍裝備部駐西安地區(qū)第四代表室,陜西 西安 710054)
近年來(lái),隨著人工智能和大數(shù)據(jù)的廣泛應(yīng)用和快速發(fā)展,高性能并行計(jì)算變得越來(lái)越重要.因此,傳統(tǒng)的從單核到多核的處理器方案已經(jīng)不能夠解決當(dāng)前計(jì)算問(wèn)題的需求.同時(shí),傳統(tǒng)處理器的發(fā)展遭遇瓶頸,無(wú)法簡(jiǎn)單通過(guò)增加核數(shù)繼續(xù)帶來(lái)運(yùn)算速度上的顯著提升.在這種情況下,出現(xiàn)了異構(gòu)硬件系統(tǒng),即不同設(shè)備或是處理器共同完成同一任務(wù).因此,傳統(tǒng)編程模型已經(jīng)不能滿足異構(gòu)硬件系統(tǒng)的并行計(jì)算需求,這直接推動(dòng)異構(gòu)并行編程模型的研究和發(fā)展.異構(gòu)并行編程模型主要包括OpenCL編程模型和CUDA編程模型.
異構(gòu)并行編程平臺(tái)能夠連通異構(gòu)系統(tǒng)和上層應(yīng)用,將復(fù)雜的異構(gòu)計(jì)算硬件平臺(tái)結(jié)構(gòu)抽象化,并形成編程接口.與串行編程模型相比,并行編程模型能夠利用異構(gòu)系統(tǒng)各設(shè)備并行計(jì)算的處理能力,以滿足當(dāng)前社會(huì)對(duì)人工智能計(jì)算及大數(shù)據(jù)計(jì)算需求.因此,異構(gòu)并行編程模型和異構(gòu)硬件平臺(tái)的研究成為當(dāng)今計(jì)算技術(shù)發(fā)展的重要組成部分.
異構(gòu)并行編程面臨兩個(gè)問(wèn)題:第一,任務(wù)分配問(wèn)題,這是由于異構(gòu)系統(tǒng)中設(shè)備的運(yùn)算速度不同所導(dǎo)致.第二,內(nèi)存管理問(wèn)題,這是由于異構(gòu)設(shè)備中計(jì)算單元內(nèi)存分配方式不同所引起.
目前NVIDIA提出的CUDA異構(gòu)并行編程模型,與Khronos提出的OpenCL編程模型等解決方案,都實(shí)現(xiàn)了異構(gòu)并行編程框架的設(shè)計(jì)思想.本文選用這兩種典型的編程模型對(duì)異構(gòu)并行編程模型進(jìn)行說(shuō)明,通過(guò)比較分析異構(gòu)并行編程模型的主要支持機(jī)制,并且展望異構(gòu)并行編程模型發(fā)展趨勢(shì).
NVIDIA 在2007 年提出了對(duì)應(yīng)NVIDIA GPU 的并行編程模型CUDA.初始,NVIDIA 將CUDA 的出現(xiàn),作為GPU并行編程進(jìn)入主流編程的一個(gè)契機(jī),就是通過(guò)CUDA編程模型為GPU增加一個(gè)可以使用的編程接口.無(wú)論使用NVIDIA生產(chǎn)的何種GPU相關(guān)設(shè)備,編程人員都可以通過(guò)CUDA進(jìn)行GPU編程.所以,CUDA最初定義為統(tǒng)一計(jì)算架構(gòu)(Compute Unified Device Architecture).編程人員無(wú)須學(xué)習(xí)復(fù)雜的著色語(yǔ)言或者圖形處理原理,也能夠進(jìn)行GPU編程[1].
OpenCL 最初由蘋果公司設(shè)計(jì)并研發(fā),之后交給開(kāi)源組織平臺(tái)Khronos Group 進(jìn)行維護(hù).作為異構(gòu)并行編程的一種平臺(tái)標(biāo)準(zhǔn),幫助編程人員能夠在異構(gòu)硬件系統(tǒng)中編寫程序.OpenCL 與CUDA 不同點(diǎn)在于OpenCL的API可以針對(duì)不同的生產(chǎn)商的設(shè)備(如:GPU、DPS、FPGA等)進(jìn)行編程[2]
編程模型的出現(xiàn)與發(fā)展離不開(kāi)硬件技術(shù)的進(jìn)步.由于CPU運(yùn)算能力十分強(qiáng)大,傳統(tǒng)的串行編程模型一直在計(jì)算機(jī)編程中占據(jù)主流地位,從面向過(guò)程的C語(yǔ)言到面向?qū)ο蟮腃++與JAVA,都是串行編程中得到廣泛應(yīng)用的編程模型.同樣,異構(gòu)并行編程模型的興起離不開(kāi)GPU的發(fā)展,因?yàn)镚PU不僅保留了CPU的流水線處理能力,并且能夠?qū)崿F(xiàn)SIMD(單指令多數(shù)據(jù))或者M(jìn)IMD(多指令多數(shù)據(jù))的執(zhí)行方式.其中,NVIDIA 提出自己的GPU 硬件調(diào)度方式,即SPMD(單程序多數(shù)據(jù)).通過(guò)SPMD 的形式執(zhí)行數(shù)據(jù)并行性,CUDA能夠?qū)崿F(xiàn)線程、線程束到線程塊的線程網(wǎng)格結(jié)構(gòu),在異構(gòu)并行編程模型的發(fā)展起到重要作用.
現(xiàn)階段,由于異構(gòu)硬件系統(tǒng)的本質(zhì)特點(diǎn),即系統(tǒng)是由不同設(shè)備或者處理器組成(如CPU+GPU、CPU+DPS、CPU+FPGA 等),異構(gòu)并行編程模型需要依靠統(tǒng)一的編程接口完成任務(wù)分配與內(nèi)存管理等問(wèn)題.程序運(yùn)行在主機(jī)端需要串行編程,并通過(guò)統(tǒng)一的編程接口,將任務(wù)分配給并行運(yùn)算的設(shè)備端,進(jìn)行并行編程.因此,異構(gòu)并行編程模型與傳統(tǒng)的串行編程模型相比,區(qū)別在于是否通過(guò)統(tǒng)一的異構(gòu)并行接口完成所需要的并行運(yùn)算,如圖1所示.
在異構(gòu)并行編程模型中,CUDA和OpenCL主要通過(guò)對(duì)于C語(yǔ)言進(jìn)行擴(kuò)展,實(shí)現(xiàn)異構(gòu)并行編程接口的技術(shù)方案.因此,通過(guò)CUDA編程模型與OpenCL編程模型分別和C語(yǔ)言進(jìn)行對(duì)比,說(shuō)明異構(gòu)并行編程模型與傳統(tǒng)串行編程模型的區(qū)別.
CUDA編程模型的異構(gòu)并行編程接口的形式是對(duì)C語(yǔ)言進(jìn)行異構(gòu)并行擴(kuò)展,通過(guò)CUDA Library的方式擴(kuò)展,可以解決任務(wù)分配與內(nèi)存分配的問(wèn)題.CUDA將串行部分與并行部分分別定義為主機(jī)端(host)與設(shè)備端(device).主機(jī)端與設(shè)備端之間相互連接的并行連接點(diǎn)是內(nèi)核函數(shù).CUDA矩陣乘法的偽代碼為:
矩陣乘法,是用矩陣A每行與矩陣B的每列依次乘積累加,如此得到各個(gè)元素的值.在傳統(tǒng)串行編程模型中,使用三層循環(huán).在異構(gòu)并行編程模型中,CUDA 使用共享內(nèi)存方法,將C(i,j)=sum{A(i,k)*B(k,j)}再度細(xì)分.CUDA將N*N的矩陣再度劃分成n*n的子塊,然后每個(gè)block負(fù)責(zé)計(jì)算子塊i和子塊j的子乘積,計(jì)算完后將其相加.
而在設(shè)備端,內(nèi)核函數(shù)引入了threadIdx、blockIdx與blockDim的變量.這些變量提供了一種跨域元素(如向量、矩陣或卷)調(diào)用計(jì)算的途徑.線程的索引和線程ID以一種直觀的方式相互關(guān)聯(lián):對(duì)于一維向量,它們是相同的;對(duì)于二維向量大小為(Dx,Dy),與之對(duì)應(yīng)的索引(x,y)的線程ID為(x+y*Dx);對(duì)于大小為(Dx,Dy,Dz)的三維向量,索引線(x,y,z)的線程ID是(x+y*Dx+z*Dx*Dy).
與CUDA相同,OpenCL編程模型也是通過(guò)對(duì)C語(yǔ)言進(jìn)行異構(gòu)并行擴(kuò)展,解決任務(wù)與內(nèi)存的技術(shù)難點(diǎn).OpenCL也是將程序分為主機(jī)端(host)與設(shè)備端(device).
在OpenCL編程模型中,最重要的部分就是設(shè)備端的內(nèi)核函數(shù)Kernel_function(),并且,矩陣運(yùn)算最能體現(xiàn)OpenCL價(jià)值的地方.在N*N的矩陣乘法中,所有乘法部分都可以進(jìn)行并行.N*N的矩陣相乘,那么進(jìn)行N*N*N 次的乘法.在異構(gòu)并行硬件平臺(tái)內(nèi),可以進(jìn)行N*N*N 次的并行運(yùn)算.OpenCL 矩陣相乘的偽代碼為:
在OpenCL中,內(nèi)核函數(shù)聲明使用kernel或__kernel修飾符來(lái)通知編譯器,這是一個(gè)OpenCL C內(nèi)核函數(shù).內(nèi)核只包含計(jì)算的代碼,在數(shù)據(jù)并行模式下,同時(shí)會(huì)有多個(gè)工作項(xiàng)(異構(gòu)并行硬件平臺(tái)的計(jì)算單元)來(lái)參與計(jì)算,每個(gè)工作項(xiàng)只在自己ID有關(guān)的計(jì)算中運(yùn)行.在代碼中,使用內(nèi)部函數(shù)get_global_id()來(lái)獲得當(dāng)前工作項(xiàng)的ID.工作項(xiàng)雖然執(zhí)行相同的代碼,但是通過(guò)ID來(lái)確定每個(gè)工作項(xiàng)的任務(wù)數(shù)據(jù),由此實(shí)現(xiàn)SIMT模式的執(zhí)行流程.
在傳統(tǒng)串行編程中,串行編程接口不需要考慮多種設(shè)備與設(shè)備并行的問(wèn)題.在面對(duì)異構(gòu)硬件系統(tǒng)與并行程序運(yùn)行時(shí),異構(gòu)并行編程需要提供不同設(shè)備的并行任務(wù)分配機(jī)制.因此,異構(gòu)硬件系統(tǒng)中的設(shè)備描述以及并行程序中的任務(wù)并行是異構(gòu)并行編程模型的兩個(gè)重要方面.CUDA與OpenCL都通過(guò)C語(yǔ)言的擴(kuò)展方案,實(shí)現(xiàn)了以上兩個(gè)方面.但在并行性任務(wù)分配上,它們都采取了與線程相關(guān)的SPMD(單程序多數(shù)據(jù)).下面將討論與研究CUDA與OpenCL的并行任務(wù)分配方式.
NVIDIA的硬件調(diào)度方式為SPMD(單程序多數(shù)據(jù)),它基于NVIDIA硬件層的支持,屬于SIMD(單指令多數(shù)據(jù))指令方式.通過(guò)SPMD的硬件調(diào)度方式,CUDA能夠?qū)崿F(xiàn)線程、線程束再到線程塊、線程網(wǎng)格的結(jié)合,和對(duì)設(shè)備的抽象化、可視化的管理.
線程是CUDA并行程序的基本構(gòu)成.在線程模型中,線程是編程架構(gòu)的最基本邏輯單位,可以完成一次簡(jiǎn)單的邏輯運(yùn)算.因此,每個(gè)線程都擁有單獨(dú)的尋址操作和寄存器處理.線程的獨(dú)立性特點(diǎn),有利于完成并行運(yùn)算.線程束是CUDA并行程序的基本執(zhí)行單元,其中包含一定數(shù)量的線程.CUDA程序通過(guò)線程束調(diào)用多個(gè)線程.線程塊是包含多個(gè)線程束的單位,它具有維度的屬性.線程塊的設(shè)計(jì)來(lái)自于圖像處理的圖形化單位.三個(gè)維度分別代表一維、二維和三維的線程結(jié)構(gòu),即多個(gè)線程可以組成圖形的效果.
線程網(wǎng)格是抽象的單位,是幫助程序員進(jìn)行編程的時(shí)候能夠更容易理解線程運(yùn)行的方式.同樣,線程網(wǎng)格也具有最高為3的維度.CUDA通過(guò)線程、線程束和線程塊,最終組成線程網(wǎng)格.線程網(wǎng)格能夠?qū)⑷蝿?wù)分配進(jìn)入線程中,使得異構(gòu)并行程序能夠在CUDA對(duì)應(yīng)的硬件平臺(tái)運(yùn)行.
OpenCL作為一個(gè)面向異構(gòu)系統(tǒng)的開(kāi)源編程模型,它為開(kāi)發(fā)人員提供了一個(gè)編程框架以及一個(gè)低級(jí)別的硬件抽象層來(lái)支持開(kāi)發(fā)并行程序.所以,與CUDA相同,OpenCL也將硬件進(jìn)行抽象化的描述,即平臺(tái)模型.平臺(tái)模型是由異構(gòu)系統(tǒng)中多臺(tái)設(shè)備共同組成,其中包括主機(jī)端與設(shè)備端.這里,OpenCL設(shè)備可以能是CPU、GPU、DSP或者FPGA等.OpenCL程序在編譯時(shí),分為主機(jī)端程序和設(shè)備端程序,這與CUDA是有相似之處的.OpenCL內(nèi)核是用OpenCL編程語(yǔ)言編寫,并用OpenCL編譯器編譯的函數(shù).
OpenCL程序包含主機(jī)端程序和設(shè)備端程序.主機(jī)端程序運(yùn)行在主機(jī)處理器上,主機(jī)端程序以命令方式將內(nèi)核程序從主機(jī)提交到OpenCL設(shè)備,OpenCL設(shè)備在處理單元上執(zhí)行計(jì)算.根據(jù)這兩個(gè)不同執(zhí)行單元定義了OpenCL執(zhí)行模型.
在確定平臺(tái)模型之后,OpenCL程序需要確定執(zhí)行模型,這是OpenCL程序建立過(guò)程中重要組成部分.執(zhí)行模型是OpenCL創(chuàng)建上下文,建立命令隊(duì)列,通過(guò)OpenCL平臺(tái)獲取信息,創(chuàng)建程序?qū)ο?
根據(jù)OpenCL設(shè)備編譯、構(gòu)建程序?qū)ο蟮牟煌瑒?chuàng)建內(nèi)核對(duì)象.經(jīng)歷這些操作,基本完成了OpenCL在設(shè)備內(nèi)核上的編譯.對(duì)于OpenCL設(shè)備上執(zhí)行函數(shù)所需要的參數(shù),需要設(shè)置參數(shù)函數(shù)將參數(shù)寫入在異構(gòu)系統(tǒng)中執(zhí)行工作的工作組和工作項(xiàng).當(dāng)以上所有操作完成后,通過(guò)命令隊(duì)列將函數(shù)入隊(duì),并交給相應(yīng)的OpenCL設(shè)備執(zhí)行.
異構(gòu)并行編程模型對(duì)于異構(gòu)架構(gòu)和上層應(yīng)用,起到橋梁的作用.與傳統(tǒng)的數(shù)據(jù)串行編程模型不同,異構(gòu)并行編程模型能夠更加充分利用異構(gòu)系統(tǒng)的豐富硬件資源,并且提供了人性化的接口與編譯器使編程人員能夠完成上層應(yīng)用的設(shè)計(jì)與編寫.
CUDA編程模通過(guò)自身的線程機(jī)制和存儲(chǔ)機(jī)制,解決了設(shè)備之間的運(yùn)行與通信問(wèn)題,而且通過(guò)編程接口解決了并行數(shù)據(jù)配置和計(jì)算的同步異步問(wèn)題.OpenCL則通過(guò)將設(shè)備抽象化,建立平臺(tái)模型,通過(guò)上下文機(jī)制和命令隊(duì)列,解決不同設(shè)備之間的問(wèn)題,而且,通過(guò)建立內(nèi)核對(duì)象和程序?qū)ο?,解決數(shù)據(jù)配置與存儲(chǔ)、計(jì)算的同步異步問(wèn)題.