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