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%。总的来说,并行编程模型在各平台的测试中额外的性能损耗都没有超过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 等通信技术实现在集群分布式环境的并行计算。