ParM:基于國產處理器的異構并行編程模型

朱文龍,江嘉治,黃 聃,肖 儂

(中山大學計算機學院,廣東 廣州 510006)

1 引言

隨著人工智能、科學計算等領域算力需求的增長,異構多核設備成為高性能計算發展的熱點方向。近年來各種異構計算芯片不斷出現,ARM[1]等精簡指令集架構開始走向服務器領域。硬件設備的多樣化發展帶來了越來越分裂的編程生態,例如 NVIDIA GPU專用的CUDA編程模型[2]、 AMD 領導開源的ROCm平臺和HIP (Heterogeneous-computing Interface for Portability)編程模型[3]、Apple芯片使用的Metal編程模型[4]和華為昇騰AI處理器使用的AscendCL編程模型[5]。此外,還有用于CPU 并行的編程模型,例如OpenMP[6]和 Pthread[7]。隨著芯片國產化浪潮的到來,國內新興的編程模型還會越來越多。

編程模型的多樣化給并行應用的開發人員帶來了很大的負擔,開發人員不僅需要精通各種硬件的底層架構,還需要同時維護多個代碼庫,這給并行應用的開發和維護帶來了越來越大的成本,阻礙了并行應用的發展。目前國外已經出現了一些支持多種體系結構的統一編程模型,但是針對國產計算設備的統一編程模型的研究與實現還比較少。

基于上述問題,本文設計實現了一個性能可移植的異構并行編程模型ParM。該模型可以在多種硬件設備上運行,能夠有效地降低開發和維護并行應用的成本。本文的主要工作如下:

(1)設計實現了一個統一的異構并行編程模型ParM,通過內存管理模塊實現了對各種異構設備上的內存抽象;通過并行計算模塊實現了對異構設備上并行操作的抽象。該模型可以在 x86 CPU、NVIDIA GPU、ARM CPU和昇騰AI處理器上高效運行。

(2)針對ARM眾核架構提出了動態負載均衡、多線程池和數據劃分模型的優化策略,能有效提升ARM眾核架構上并行計算的性能。

(3)把昇騰AI處理器的使用抽象為數據模型和任務模型,用數據模型管理昇騰設備內存,用任務模型封裝昇騰算子的調用流程,簡化了昇騰AI處理器的使用步驟。

2 背景

2.1 異構并行編程模型

異構并行編程模型[8]是隨著異構計算設備發展而興起的、支持不同體系結構的統一編程框架,它對各種設備的架構進行抽象,統一成高層次編程接口,再映射到底層并行庫實現。異構編程模型的優勢在于開發者不需要關注設備的具體架構,開發應用方便,而且代碼具有可移植性,能在多種設備上運行。

國外的異構編程模型經過多年的發展,已經產生了一些應用廣泛的編程模型,例如Kokkos、RAJA和SYCL等。

Kokkos[9]是由美國Sandia國家實驗室領導開發的,目前已經發展出了一套完整的生態系統,在科學計算領域得到了廣泛使用。Kokkos是一個C++庫,通過模板元編程[10]的方式把Kokkos代碼在編譯期轉換為指定設備的底層代碼,從而實現了性能可移植性。目前Kokkos支持的并行后端有CUDA、OpenMP、Pthread、HIP和SYCL。

RAJA[11]是美國勞倫斯利弗莫爾國家實驗室LLNL(Lawrence Livermore National Laboratory)開發的一個面向高性能計算應用程序的C++抽象軟件庫,提供了架構和編程模型的可移植性。RAJA為循環提供了可移植的抽象,將應用程序循環內核與底層架構和特定于編程模型的實現細節隔離開來。支持多種執行后端,具備性能可移植性。

OpenCL(Open Compute Library)[12]是一個基于數據分割和任務分割并行計算機制的異構編程模型,支持CPU、GPU 和FPGA等不同設備,具備極強的通用性,但其程序繁瑣,開發復雜。Khronos發布了一個基于OpenCL的高層編程模型SYCL[13],同樣基于C++特性,SYCL通過上層接口大大簡化了OpenCL的接口調用,支持模板和Lambda函數技術,大大增強了編程的靈活性。

目前國內對異構并行編程模型的相關研究還比較少,有基于OpenCL實現的UPPA(Unified Parallel Programming Architecture)[14]等。隨著國產計算設備的發展[15],針對國產計算設備的異構編程模型將是未來研究的一個重要方向。

2.2 華為計算設備

本文設計的異構并行編程模型重點面向華為計算設備,包括華為鯤鵬處理器和華為昇騰 AI 處理器。

華為鯤鵬處理器[16]是華為公司基于ARM指令集自主研發的非統一內存訪問NUMA(Non- Uniform Memory Access)架構處理器,包括多個 NUMA結點,每個結點對應一部分內存。鯤鵬處理器中的NUMA架構如圖1所示。

Figure 1 Architecture of NUMA

華為昇騰AI處理器是華為公司研發的AI處理器,主要面向AI計算任務。昇騰AI處理器包括AI CPU和AI Core 2個部分,AI CPU主要用于控制算子執行等通用計算;AI Core是昇騰AI處理器的計算核心,負責執行與向量和張量相關的計算密集型算子。AI Core使用了華為自主研發的達芬奇架構[17],每個AI Core中實現了3種計算單元:矩陣計算單元、向量計算單元和標量計算單元。其中,專為矩陣乘法實現的矩陣計算單元能在1個時鐘周期內完成16×16的矩陣乘法,有效提升了深度學習中矩陣運算的效率。

昇騰AI處理器使用AscendCL作為開發框架,AscendCL以算子為基本執行單位,程序通過調用算子完成特定的并行操作。AscendCL中內置了一些深度學習和圖像處理常用的算子,用戶也可以在AscendCL中注冊自定義算子。

3 ParM整體架構

ParM以C++庫的形式提供,其設計思想是用C++模板元編程的方式在編譯期將ParM代碼轉換為對應底層設備的代碼。ParM整體架構如圖2所示,其編程模型分為內存管理模塊和并行計算模塊2個部分。在內存管理模塊中,使用數據視圖作為內存在上層的接口,用引用計數的方式對內存進行管理,通過內存分配器把數據視圖映射到不同的底層設備內存。在并行計算模塊中,通過執行空間、計算策略和計算模式的不同組合來表示各種并行操作,并把并行操作映射到后端設備上。

Figure 2 Overall architecture of ParM

3.1 內存管理模塊

內存管理模塊實現了通用內存表示在底層硬件上的映射。主要包括2個部分:內存分配器 VDATA_Allocator和數據視圖VDATA_VIEW。本文設計的內存管理模塊支持4種底層內存:HostMemory、CUDAMemory、NUMAMemory和AscendMemory,分別表示x86 CPU、NVIDIA GPU、鯤鵬CPU和昇騰AI處理器上的內存空間。

內存分配器VDATA_Allocator實現了在不同內存空間的內存分配和釋放,以及數據在主機端和設備端的相互傳輸。內存分配器接收2個模板的參數:數據類型和內存空間,然后把對內存的操作映射到對應后端的內存操作。例如,內存分配操作分別映射到后端的malloc、cudaMalloc、numa_alloc_onnode 及aclrtMalloc。

數據視圖是基于內存分配器實現的具有通用性和性能可移植性的多維數據存儲。數據視圖 VDATA_VIEW最多有8個維度,通過內存分配器映射到底層的一維數據結構中。創建數據視圖需要指定的參數有:數據類型(例如int、float 等)、存儲空間(如CUDAMemory、AscendMemory等)、映射方式(如行優先、列優先等)和訪問屬性(如原子操作)。通過數據視圖VDATA_VIEW的參數組合來表示設備上的內存空間。數據視圖內部使用引用計數的方式對內存進行管理,同時使用顯式深拷貝函數進行內存復制。下面展示了數據視圖聲明示例:

using type=float;

//聲明數據視圖類型view_type

using view_type=ParM::VData〈type**〉;

//聲明二維數據視圖x

view_typex("x",10,10);

3.2 并行計算模塊

并行計算操作在ParM框架中被抽象為執行空間、計算模式和計算策略的組合。其中,執行空間定義了計算使用的后端框架,計算策略表示任務分配的方式,計算模式表示具體執行的并行操作。并行操作的內容用C++仿函數或λ表達式表示,使用執行空間、計算策略和計算模式3個抽象來描述并行計算在哪種設備以怎樣的方式執行。三者的不同類型可以相互組合,完成各類不同類型的并行計算任務,并行編程模型通過并行計算抽象可以靈活地將程序的并行計算任務映射到不同體系結構的計算設備上。

3.2.1 計算模式

計算模式表示具體執行的并行操作。本文定義了3種計算模式:簡單并行(Fork-join)、映射(Mapping)和模板(Stencil),分別對應于提供的3個并行編程接口:parallel_exec()、parallel_cal()和parallel_stencil()。3種計算模式示意如圖3所示。

Figure 3 Execution patterns

3個并行編程接口的調用方式如下所示:

//簡單并行模式接口使用

parallel_exec〈OpenMPSpace〉(8,(){…});

//映射模式接口使用

parallel_cal〈OpenMPSpace〉(Policy,(inti){…});

//模板模式接口使用

parallel_stencil〈OpenMPSpace〉(Policy,StencilFunctor);

Fork-join模式下模板參數為執行空間,傳入的第1個參數為線程數,第2個參數為表示執行任務的仿函數或Lambda表達式。Fork-join模式直接把任務分配到指定的線程上,因此不需要指定計算策略,是最簡單的并行計算模式。

Mapping模式下傳入的第1個參數為計算策略,第2個參數為表示執行任務的仿函數或Lambda表達式。在Mapping模式下任務根據傳入的計算策略分配到對應的執行空間執行,例如使用區間策略可以實現單層循環并行,使用多維區間策略可以實現多層循環并行,使用分組策略可以實現分組并行。

Stenci模式表示模板運算,傳入參數為 StencilFunctor。用戶在StencilFunctor中定義如何通過鄰居元素和當前元素計算出一個新的值,并填入到輸出集合。例如,2d-9pt模板計算表示二維數據每次計算都會使用當前計算位置和周圍8個鄰居的數據(9個值)。

3.2.2 計算策略

ParM使用計算策略來表示任務執行的方式,分為區間策略(RangePolicy)、多維區間策略(MDRangePolicy)和分組策略(GroupPolicy)。下面展示了3種策略的聲明方式:

//區間策略

autopolicy1=newRangePolicy(0,100);

//多維區間策略

autopolicy2=newMDRangePolicy({{0,0},{100,100}});

//分組策略

autopolicy3=newGroupPolicy(2,8);

區間策略用一個區間范圍來表示一維連續執行的任務索引集(如[0,H]),區間范圍內的每個元素都僅被執行一次,元素之間沒有執行順序上的依賴,在該策略下每一個工作項都是相互獨立進行并行計算的。

多維區間策略把任務劃分到一個多維空間中,空間中每個任務相互獨立且僅執行一次,用于需要對多維數組操作的情形。在多維區間策略中,支持用戶對任務進行分塊,并行后端會根據分塊大小優化數據訪問。例如,CUDA執行空間中會根據分塊大小在線程塊(block)中設置相應的線程數,根據分塊數目設置block數。用戶可以根據不同的并行后端進行合理分塊,以充分利用硬件性能。

分組策略表示把任務劃分到線程組執行,用于需要分層并行的情況。例如,定義2個線程組,每組8個線程。任務分為2級并行,組間并行和組內并行。線程組根據執行空間的不同進行映射,例如在NUMA架構CPU上同一個線程組映射到同一個NUMA結點,在GPU上一個線程組映射到一個block中。使用分組策略時,仿函數和Lambda執行函數的參數不再是執行索引,而是代表分組成員信息的句柄GroupMember。該句柄中包括分組ID、分組大小和分組中的線程ID,可以用于實現分組同步功能。分組策略需要2個參數:分組數量和分組大小。分組數量代表外層迭代范圍,數量上不設限制,但其并行性受到硬件條件約束;分組大小則會受制于硬件條件,通過啟發式方式得到最大分組線程數量,在CPU 環境由計算核心數量確定,GPU環境上則由寄存器、內存、流式多處理器SM(Streaming Multiprocessor)數量等硬件屬性決定。在分組并行中使用線程屏障(Barrier)自主控制線程行為,避免在進行內存循環時出現不同步現象。CPU類后端的線程屏障基于內存計數實現,CUDA后端則是直接使用同步函數__syncthreads()實現。

3.2.3 執行空間

執行空間代表各種不同計算設備所支持的編程模型后端,如面向CPU的OpenMP、Pthread 和Serial,面向GPU的CUDA以及面向晟騰AI處理器的Ascend。對于不同的執行空間,并行編程模型會將計算模式和計算策略定義的執行方式映射到底層不同體系結構的計算設備上執行。執行空間承載了對各個后端編程接口的調用,包括資源初始化、去初始化等設備資源管理,以及對計算任務的并行調用執行邏輯。為了提高運行時性能,本文使用C++模板元編程技術,在編譯期就對執行空間完成實例化,以減少運行時的類型選擇。

OpenMP和Pthread執行空間是面向x86處理器和國產ARM處理器的并行后端,在并行計算任務管理上采用相同的方式,不同之處在于并行編程模型對Pthread線程需要進行創建和管理。為了對線程已分配和正在執行的計算任務增強管理,并行編程模型為OpenMP執行空間和 Pthread執行空間中的每一個線程維護了一份私有數據ExecPartition,用于記錄線程管理信息和任務信息。OpenMPSpace和PthreadSpace在初始化時就提前啟動線程,當正式執行并行操作時就可以更快響應,用戶可以自定義配置線程綁定到特定的計算核心。直到去初始化時才終止線程,釋放私有數據空間。在OpenMP執行空間中,本文只使用了OpenMP的多線程管理能力,即只使用了OpenMP的編譯指令#pragma omp parallel,將多線程的任務管理保留在并行編程模型,通過線程管理和任務管理分離的方式,使得任務的分配和管理變得更加靈活,為進一步優化提供了空間。

CUDA 執行空間是面向GPU的并行后端,CUDA執行空間在初始化時,會獲取GPU設備的屬性參數(流式多處理器(SM)信息、線程塊最大限制、線程束數量限制等等)來初始化CUDA后端執行環境,以便后續管理并行計算的線程規模。CUDA編程模型的并行任務需要使用核函數啟動執行,并傳入線程塊和線程塊網格的大小。本文CUDA執行空間中實現了一個統一的核函數,由統一核函數調用執行體函數。CUDA后端的并行程度由線程塊和線程塊網格大小決定,合理設置線程塊和線程塊網格可以匹配GPU存儲架構,提高并行計算的性能,CUDA執行空間根據不同的計算策略提供了不同的線程塊和網格配置方法,使得計算任務可以更好地并行執行。

Ascend執行空間是面向昇騰AI處理器的并行后端,通過使用多個任務流(Stream)來實現計算任務的并行。在主機端使用單線程多Stream的方式異步啟動多個算子并行執行,然后進入同步等待狀態,直到計算完成。計算過程中Stream的數量由用戶指定,默認為1,使用Range計算策略表示任務總數,把任務平均劃分到各個Stream中的算子。計算任務以算子的形式提交,在昇騰設備端的各個計算核心(AI Core 或 AI CPU)上進行并行處理。由于算子之間不可進行同步操作,所以Ascend執行空間只支持不存在相互依賴的數據并行。AscendSpace在使用前需要先進行AscendCL初始化,才能使用設備資源,然后需要申請運行資源,包括綁定昇騰設備、創建運行時上下文及創建并行執行需要使用到的多個Stream。這些操作都在Ascend執行空間初始化時自動完成,執行資源在執行空間中維護,直到并行計算任務完成后去初始化。Ascend后端架構如圖4所示。

Figure 4 Backend implementation of Ascend parallel

4 ParM性能優化

由于統一編程模型接口增加了異構設備體系結構相關的處理邏輯和并行計算抽象解析,性能損失難以避免。為進一步提高并行編程模型的性能可移植性,本節介紹了針對異構設備不同的架構特征進行的性能可移植性優化。

4.1 基于均衡分配和動態負載的任務分配策略

本文在分配任務時使用均衡的任務分配策略,即把任務平均分給每個線程。然而,由于問題本身的稀疏性和不規則性,靜態的分配策略往往不能保證任務在運行過程中始終保持負載均衡。針對可能存在的負載不均衡問題,本文在ParM編程模型中使用了工作竊取[18](Work-Stealing)策略。在工作負載均衡策略中,并行編程模型會為每一個線程創建一個任務隊列,用來保存分配給線程的計算任務。并行任務開始時會進行一次近似完全均衡策略的任務分配,將任務分配結果寫入到各個線程的雙端任務隊列中。當某個線程計算完任務隊列中的任務時,線程會嘗試從下一個線程的任務隊列中提取最后一個任務進行計算。如果當前線程的下一個線程隊列也為空時,會沿著并行線程組成的線程環依次檢測對應的任務隊列的最后一個線程執行,直到再次回到線程自己本身的任務隊列時,線程退出本次并行。算法1展示了工作竊取策略的調用步驟。

算法1工作竊取算法

輸入：待計算任務集合S,任務數量m,線程數n。

輸出：計算結果。

Step1初始化線程集合T。

Step2初始化雙端隊列集合Q。

Step3將m個任務均勻分配到n個雙端隊列,每個雙端隊列中有m/n個任務。

Step4用工作竊取算法計算任務。

foreach thread inTdo

whileQ[thread_id]!=empty()do

計算Q[thread_id]隊頭任務;

Q[thread_id].pop_front();

end

fori→thread_idtondo

ifQ[i]!=empty()then

計算Q[i]隊尾任務;

Q[i].pop_end();

end

end

forj→ 0 tothread_id-1do

ifQ[j]!=empty()then

計算Q[j]隊尾任務;

Q[j].pop_end();

end

end

end

4.2 NUMA感知的多線程池和數據劃分模型

在 NUMA共享內存體系結構中,處理器直接訪問本地NUMA結點的內存時訪問速度較快,使用共享總線訪問屬于其他 NUMA結點的內存時訪問速度較慢。處理器的親和性和數據的內存位置是影響NUMA架構程序最重要的2個方面,特別是在ARM NUMA 體系架構下,NUMA內存訪問的這個特性表現得更加明顯。針對NUMA共享內存體系結構的訪問特點,本文提出了多線程池和數據劃分模型,以提升編程模型在鯤鵬處理器上的性能。

為了減少跨NUMA結點的內存訪問,本文將線程綁定到特定NUMA結點上,線程內處理和生成的數據均盡可能保存在本地NUMA結點。本文首先根據NUMA結點的數量建立起等量的線程池,每個線程池綁定到一個NUMA結點上,同時設置該線程池中的線程所需的內存優先分配到當前結點。在編程模型分配任務時首先用靜態分配的方式在線程池之間均勻分配任務,然后再把任務分配到線程池內的各個線程,池內的線程之間支持動態負載均衡。

在NUMA架構上,C/C++原始的內存分配接口malloc()會根據First-Touch內存分配原理[19],在線程所在的本地NUMA結點開辟物理內存存放數據。這種方式會造成多個NUMA結點上的線程必須跨結點訪問數據內存,不利于并行訪問。

針對這一問題,本文設計了一個多NUMA結點的數據分配接口numa_alloc(),將申請的內存平均分配到多個NUMA結點,使得后續的并行數據處理過程可以盡可能使用本地NUMA 結點內存上的數據。該接口的內存分配由數據劃分模型實現,如圖5所示,在數據劃分模型中保存了指向每個 NUMA結點的指針。內存分配實現了2種方式:第1種方式是將整個內存進行劃分,把每一塊分配到一個NUMA結點,分塊的大小可以用戶自定義,默認是平均分塊;第2種方式是在每個 NUMA結點上申請同樣大小的內存空間,這種分配方式一般用于在各個NUMA結點內存上保存相同的數據。數據劃分模型重載了操作運算符“[]”,方便用戶像使用C/C++原有的申請內存的指針一樣使用數據劃分模型。

Figure 5 Data partition model

使用數據劃分模型將數據分布到多個 NUMA結點,后續使用多線程池的線程管理方法并行時,各個NUMA結點的線程池中的線程會根據本地NUMA編號,在數據劃分模型中得到分配在當前結點的數據指針,從而實現本地數據訪問。多線程池和數據劃分模型結合的線程和內存管理方法對于解決無依賴的數據并行和局部有數據依賴(如模板計算)的并行計算任務有較大的優勢,可以最大化本地數據的訪問,減少跨NUMA結點的內存訪問。

4.3 基于多層次并行的模板計算

在NUMA架構的并行計算中可能會出現跨 NUMA結點訪問,例如模板計算中會用到鄰居的數據。本文使用向量指令并行、線程并行和 NUMA結點并行三級并行的方式優化跨 NUMA結點計算的數據訪問。

本文根據NUMA結點數量對模板計算的輸入數據進行分割后放置到各個NUMA結點,在劃分時針對輸入數據的最外層進行劃分,例如 2D數據沿著y軸劃分,3D數據沿著z軸劃分。圖6展示了2D數據的劃分方式,輸入數據被平均劃分為4份,分別存放到4個NUMA結點上。其中,只有內部邊界數據在每次迭代時需要在 NUMA結點之間進行交換,這類數據稱為光環數據。

Figure 6 2D data segmentation

在數據劃分之后,使用多線程池方法對各個 NUMA結點的數據進行計算,各NUMA結點間相互獨立,只需在每次迭代后進行通信和交換光環數據,這樣就能保證下一次迭代使用到新的數據,并且可以減少通信的數據量。在線程內部,本文使用NEON指令進行SIMD并行,每個向量寄存器長度為128位,可以同時進行2個雙精度浮點數或4個單精度浮點數運算。

4.4 昇騰AI處理器調用流程優化

由于昇騰AI處理器以算子為單位調用執行,且主要面向AI計算任務,因此本文沒有把昇騰加速卡作為通用計算設備使用,而是把昇騰算子的調用流程集成到本文設計的編程模型中,通過編程模型對昇騰加速卡的自定義算子調用流程進行簡化。

本文把昇騰自定義算子的調用抽象成數據模型和任務模型2個部分。數據模型用于管理主存和昇騰設備內存,實現了昇騰設備上的內存申請以及數據傳輸,可以自動生成描述數據的 json字符串。數據模型集成到數據視圖中,用數據視圖統一表示。任務模型在parallel_cal()接口內部使用,傳入參數為算子名和數據類型,在運行時動態生成昇騰張量編譯器ATC(Ascend Tensor Compiler)工具需要的json文件,并使用基于系統命令的運行時命令調用ATC工具生成算子模型文件,準備完成后在Stream中執行算子。如圖7所示,調用流程簡化后,并行編程模型的用戶在使用時無需過多關注與并行無關的操作,只需要準備數據和調用并行接口,可以更關注于并行計算本身的邏輯。

Figure 7 Optimization of operator call process

在進行了自定義算子開發和調用流程的抽象簡化后,Ascend后端的使用方式跟其他后端基本一致。但是,由于昇騰本身編程模式的限制,跟GPU相比,昇騰的并行步驟繁多,會導致一定的性能損失,因此昇騰后端主要用來執行工作量較大的任務。

5 實驗

為了測試并行編程模型的性能可移植性和優化研究的效果,本文分別使用稠密矩陣乘法和直方圖統計算法。測試在多個平臺上進行,詳細測試平臺信息如表1和表2所示。

Table 1 CPU test platform

Table 2 GPU and Ascend test platforms

5.1 性能可移植性實驗

性能可移植性測試用于測試同一算法通過編程模型在不同平臺上運行時的性能。本文選取了稠密矩陣乘法作為測試樣例,與原始代碼進行比較。

在本文實驗中,稠密矩陣乘法使用多維區間策略和parallel_call()接口進行計算。多維區間策略的范圍是2層0到L,L是方陣大小。矩陣使用行優先的方式存儲,用C語言內置的隨機數函數生成浮點數初始化輸入矩陣。矩陣乘法用最簡單的3層循環方式實現,以最內層的循環作為并行域,外部2層循環用來索引任務單元分配給各個線程執行,以2層嵌套循環并行的方式執行。

下面是使用ParM編程模型在CPU設備實現簡單矩陣乘法的代碼,執行體使用仿函數實現,指定其他設備需要在VData聲明時更改參數ParM::HostMemory。

//聲明數據

using type=float;

using view_type=

ParM::VData〈type**,ParM::HostMemory〉;

view_typex("x",M,P);

view_typey("y",P,N);

view_typez("z",M,N);

//初始化環境

ParM::ParMInitArgumentsarguments;

arguments.threads_num=threadNum;

ParM::parm_init(arguments);

//定義執行體

template 〈typename Type〉

classfunctor{

view_typex;

view_typey;

view_typez;

intm,p,n;

public:

voidoperator()(inti,intj) const{

typeresult=0.0;

for(intk=0;k

result+=x(i,k)*y(k,j);

}

z(i,j)=result;

}

functor(const intm,const intp,const intn,view_typex,view_typey,view_typez)

:x(x),y(y),z(z),m(m),p(p),n(n) {};

};

//定義執行策略

using policy_type=

ParM::MDRangeInterval〈ParM::Rank〈2〉〉;

policy_typepolicy({{0,0}},{{M,N}});

//執行運算

ParM::parallel_cal(policy,functor〈type〉(M,P,N,x,y,z));

ParM::parm_finalize();

由于昇騰設備需要進行算子調用,所以在執行時需要額外進行以下數據描述和數據拷貝:

// 數據描述

ParM::DataDescdim_desc(std::vector〈int64_t〉{3},ParM::type2string〈int64_t〉());

ParM::DataDescx_desc(std::vector〈int64_t〉{x_size},ParM::type2string〈float〉());

ParM::DataDescy_desc(std::vector〈int64_t〉{y_size},ParM::type2string〈float〉());

ParM::DataDescz_desc(std::vector〈int64_t〉{z_size},ParM::type2string〈float〉());

//創建數據模型

ParM::Modelmodel;

model.input_desc.push_back(&dim_desc);

model.input_desc.push_back(&x_desc);

model.input_desc.push_back(&y_desc);

model.output_desc.push_back(&z_desc);

// 設備數據拷貝

ParM::Viewdim_view(3 *sizeof(int64_t),

ParM::AscendMemSpace::DeviceSpace);

ParM::Viewx_view(x_size*sizeof(type),

ParM::AscendMemSpace::DeviceSpace);

ParM::Viewy_view(y_size*sizeof(type),

ParM::AscendMemSpace::DeviceSpace);

ParM::Viewz_view(z_size*sizeof(type),

ParM::AscendMemSpace::DeviceSpace);

ParM::View::copyHost2Device(dim_view.ptr,

dim,3 *sizeof(int64_t));

ParM::View::copyHost2Device(x_view.ptr,

(void*)x,x_size*sizeof(type));

ParM::View::copyHost2Device(y_view.ptr,

(void*)y,y_size*sizeof(type));

ParM::View::copyHost2Device(z_view.ptr,

(void*)z,z_size*sizeof(type));

std::vector〈ParM::View *〉inputs;

inputs.push_back(&dim_view);

inputs.push_back(&x_view);

inputs.push_back(&y_view);

std::vector〈ParM::View *〉outputs;

outputs.push_back(&z_view);

//執行計算

ParM::RangeInterval〈〉policy(0,z_size);

ParM::parallel_cal(policy,"FunctorKernel",model,inputs,outputs);

實驗結果如圖8所示。可以看到,在大部分情況下使用本文模型實現的算法相比原始算法要慢,但是誤差保持在5%以內,說明本文實現的并行編程模型具有性能可移植性。

Figure 8 Test results of matrix multiplication

在GPU平臺上,并行編程模型程序的性能損失主要集中在將執行體拷貝給內核函數,以及對并行計算抽象的解析上。在各個不同數據規模下,CUDA后端實現的程序性能損失在1%左右。在昇騰平臺上,并行編程模型實現的程序和原始平臺程序都需要進行大部分相同的并行準備操作,兩者的執行性能相差不超過1%?？偟膩碚f,并行編程模型在各平臺的測試中額外的性能損耗都沒有超過5%。

5.2 SYCL對比實驗

由于SYCL編譯器有多種不同的實現方式,本文以Intel公司實現的DPC++[21]為例對SYCL異構編程模型進行對比測試。DPC++不支持ARM CPU和昇騰AI處理器,因此本節僅在x86 CPU平臺和GPU平臺進行對比。對比實驗使用5.1節中的矩陣乘法,SYCL內存管理使用緩沖區機制,編譯參數使用-O3。測試結果如圖9所示。

Figure 9 Comparison results of DPC++ performance

在CPU平臺,由于DPC++使用了緩沖區機制,在矩陣規模較小時DPC++性能低于OpenMP和ParM的,而在矩陣規模較大時,由于DPC++使用了Intel TBB作為后端,并且基于LLVM進行了編譯優化,在簡單3層循環下DPC++性能要好于OpenMP和ParM的。在GPU平臺,SYCL的性能稍差于CUDA和ParM的。

5.3 NUMA架構優化實驗

該節使用直方圖統計算法[21]對4.2節提出的基于國產ARM眾核架構的線程和內存管理優化進行測試。實驗只在ARM平臺進行,對比 OpenMP原始實現和使用優化后的并行編程模型實現的程序性能差異。

本文實驗使用的實驗數據是使用 C語言內置的隨機數函數生成的在0～255的隨機整數,模擬圖像的RGB像素值,數據規模為108。直方圖統計算法使用分治思想實現,在 OpenMP原始實現的程序中,使用C語言原始內存申請接口申請模擬像素點數組的空間,并進行隨機初始化。然后將像素值平均分配給各個線程,每個線程都維護一個私有的統計結果數組private_histo,先并行地將分得的像素值統計到private_histo中,最后在主程序線程中進行合并,統計到最終結果數組histo中。基于本文并行編程模型優化后的實驗組程序使用numa_alloc()接口申請像素點數據內存空間,然后使用區間策略和映射計算模式接口Parallel_cal()實現,其中numa_alloc()接口申請的數據會分布到ARM實驗平臺的4個NUMA結點,OpenMP線程根據多線程池的方案綁定到各個NUMA結點。

實驗結果如圖10所示?？梢钥吹?在實驗測試范圍內,優化后的程序比OpenMP原始實現的程序性能更好,而且隨著線程數量的增多,優化方法的性能提升也越大,線程數量少于8時,性能提升在20%以內,當線程增長到128時,性能提升了80%左右,且優化后的程序表現出了更好的可擴展性。

Figure 10 Test results of histogram statistical algorithm

6 結束語

本文設計和實現了一個通用的異構并行編程模型 ParM,支持 x86 CPU、ARM CPU、NVIDIA GPU 和昇騰 AI 處理器4種異構計算設備。ParM在異構計算設備中具備良好的性能可移植性。本文把編程模型上的內存抽象為數據類型、存儲空間、映射方式和訪問屬性的組合,用統一的數據訪問接口實現了對不同設備的內存訪問;通過把底層設備上的并行操作抽象為執行空間、計算策略和計算模式的組合,實現了底層設備的通用并行計算接口,并利用模板元編程技術把上層操作映射到底層設備上。

本文針對NUMA的架構特點對設備上的并行訪問進行了專門優化,有效提升了并行訪問效率。針對昇騰自定義算子調用流程復雜的特點,本文把昇騰調用流程集成到 ParM 編程模型中,屏蔽了與并行運算無關的部分,簡化了昇騰AI處理器的使用。實驗結果表明,本文實現的ParM異構編程模型能夠達到原始代碼90%以上的性能,表明ParM具有較好的性能可移植性。

后續的工作將主要集中在擴展更多體系結構設備的支持,包括更多的國產設備支持以及支持更多的并行模式。目前的研究主要集中于不同體系結構設備的并行計算,依賴于底層的專用編程模型,對于異構設備協同并行計算和異構設備間數據訪問缺少研究,在下一步將針對異構設備數據管理問題進行研究;而且本文的并行編程模型只面向單機器的并行,缺乏面向分布式環境的支持,后續將會基于 MPI 等通信技術實現在集群分布式環境的并行計算。