CN105279137A - 一种面向gpu并行的三对角矩阵方程求解方法 - Google Patents

Abstract

本发明公开了一种面向GPU并行的三对角矩阵方程求解方法，包括：将面向GPU并行的三对角矩阵进行矩阵块间依赖消除实现块内消元；将经过依赖消除的矩阵块间进行并行求解实现块间并行消元；利用主机内存和GPU显存的透明映射，进行主机和GPU之间的数据传输通路。本发明能够有效提高数据处理的块间并行度，保证正确性的同时，提高系统的运算效率。

Description

一种面向GPU并行的三对角矩阵方程求解方法

技术领域

本发明涉及高性能计算技术领域，尤指一种面向图形处理器(GPU，GraphicsProcessingUnit)并行的三对角矩阵方程求解方法。

背景技术

三对角方程组是线性代数方程组中结构形式较为简单的一种，属于带状方程组。其半带宽为1。在两点边值问题的有限差分方法、三次样条插值函数的计算、偏微分方程隐式差分方法中常常产生这类方程组。研究三对角线方程组的求解对数值计算相关领域(如航空数值计算)也有着重要的价值。

传统求解三对角方程组的串行算法比较经典的有Guass消去法，也被称为“追赶法”，虽然该算法已证明在串行计算机上是最优的算法，但是并不适合并行求解。三对角方程组的并行算法主要有递归倍增法(stone算法)、循环约化法、以及分裂法等。

同时，近年来GPU硬件性能和可编程性的提高，以及其在数值计算领域的广泛应用，为并行求解三对角方程组，提高计算性能带来了契机。分裂法由H.H.Wang在八十年代提出，算法贯彻了分而治之的原则，适用于并行求解。但是针对计算统一设备架构(CUDA，ComputeUnifiedDeviceArchitecture)架构，算法的几个中间步骤表现出适应性的不足。

原分裂法算法的主要步骤如下，以9乘9矩阵为例。

1.将矩阵分块，如图1所示。

2.依次消去每块左下方元素(a2，a5，a8，a3，a6，a9)和上方元素(c1，c4，c7)，块内串行，块间并行，如图2所示。

3.消去每块上方元素(c3，c6)，存在块间数据依赖，块间串行，如图3所示。

4.消去每块左侧元素(f4，f5，f6，f7，f8，f9)和右侧元素(g6，g7，c8，g3，g4，c5，g1，c2)，存在块间数据依赖，块间串行，如图4所示。

5.求解完毕，如图5所示。

在利用CUDA编程模型在GPU上进行实现时，如果按照上述原算法的描述进行求解，会发现在线程与线程之间，以及块与块之间存在着依赖关系，需要进行通信和同步操作。

根据上述算法描述和分析：在算法的前半部分，矩阵块内串行，矩阵块间并行；算法后半部分，矩阵块内并行，矩阵块间串行。前半部分的分块方法将矩阵分块的尺寸固定，之后就无法更改，而且块内的并行度低，加速效果差；后半部分块间串行无法利用GPU多线程的优势。。

发明内容

为了解决上述技术问题，本发明提供了一种面向GPU并行的三对角矩阵方程求解方法，能够有效提高数据处理的块间并行度，保证正确性的同时，提高系统的运算效率。

为了达到本发明目的，本发明提供了一种面向GPU并行的三对角矩阵方程求解方法，包括：将面向GPU并行的三对角矩阵进行矩阵块间依赖消除实现块内消元；将经过依赖消除的矩阵块间进行并行求解实现块间并行消元；利用主机内存和GPU显存的透明映射，进行主机和GPU之间的数据传输通路。

进一步地，所述进行矩阵块间依赖消除，具体为：根据矩阵数据分块之间的并行性，消除矩阵分块间的数据依赖，并将具有数据块间依赖的部分进行块内消元。

进一步地，所述进行矩阵块间依赖消除，具体为：将矩阵分块；依次消去每块左下方元素和上方元素，块内串行，块间并行；消去每块上方元素，存在块间数据依赖，块间串行；依次消去每块左侧最下元素，进行依赖数据预处理，块间串行；顺序消去每块左侧元素和右侧元素，块内串行，块间并行。

进一步地，所述进行并行求解，具体为：根据分块大小以及GPU线程数目，对将经过依赖消除的矩阵块间进行并行求解。

进一步地，所述进行主机和GPU之间的数据传输通路，具体为：利用主机内存和GPU显存的透明映射，使用主机CPU进行串行运算实现主机和GPU之间的数据传输通路。

与现有技术相比，本发明通过矩阵块间依赖消除算法，将原分裂法中具有数据块间依赖的部分进行预先消元，为后续块间并行消元打下基础，使得整个算法适用于GPU的并行处理；针对为面向GPU加速求解三对角矩阵方程的特点，选择合适的分块大小以及GPU线程数目，进行并行处理，是大幅提高计算性能的基础；利用主机内存和GPU显存的透明映射，一方面减少了数据传输量，另一方面简化了数据传输步骤，进一步挖掘了算法的整体性能。本发明充分考虑GPU具有强大并行处理能力的特点，并着眼于传统分裂法求解三对角矩阵算法的GPU适应性改进，面向GPU计算特征，采用依赖数据预消除的方法针对矩阵分块计算的并行性进行了优化，最终在三对角矩阵方程的求解速度方面，与传统的分裂法相比获得了较大的改进，从而推动了GPU在高性能数值计算领域广泛应用。

本发明的其它特征和优点将在随后的说明书中阐述，并且，部分地从说明书中变得显而易见，或者通过实施本发明而了解。本发明的目的和其他优点可通过在说明书、权利要求书以及附图中所特别指出的结构来实现和获得。

附图说明

附图用来提供对本发明技术方案的进一步理解，并且构成说明书的一部分，与本申请的实施例一起用于解释本发明的技术方案，并不构成对本发明技术方案的限制。

图1是现有技术中的分裂法矩阵分块示意图。

图2是现有技术中的分裂法块间并行消元示意图。

图3是现有技术中的分裂法块间同步消元示意图。

图4是现有技术中的分裂法块内并行、块间串行消元示意图。

图5是现有技术中的分裂法求解完毕示意图。

图6是本发明的面向GPU并行的三对角矩阵方程求解方法的流程示意图。

图7是本发明的算法初始状态示意图。

图8是本发明的算法串行依赖消去示意图。

图9是本发明的算法块间并行消元示意图。

图10是本发明的算法求解完毕示意图。

具体实施方式

为使本发明的目的、技术方案和优点更加清楚明白，下文中将结合附图对本发明的实施例进行详细说明。需要说明的是，在不冲突的情况下，本申请中的实施例及实施例中的特征可以相互任意组合。

在附图的流程图示出的步骤可以在诸如一组计算机可执行指令的计算机系统中执行。并且，虽然在流程图中示出了逻辑顺序，但是在某些情况下，可以以不同于此处的顺序执行所示出或描述的步骤。

在CUDA实现时，如果按照现有技术的方法进行求解，发现在线程与线程之间，以及块与块之间存在着依赖关系，需要进行通信和同步操作。最为重要的是最后对于f列和g列的消去线程之间有着依赖关系，必须串行进行，这对CUDA程序的性能是致命的损失。

本发明的算法设计主要考虑分裂法块内及块间可并行的特点，以及GPU编程模型CUDA单指令流多线程的特点，采用数据预处理消除块间数据依赖的方法，可有效提高数据处理的块间并行度。

如图6所示，本发明的面向GPU并行的三对角矩阵方程求解方法，包括：

步骤601，将面向GPU并行的三对角矩阵进行矩阵块间依赖消除。

在本步骤中，提出了一种消除块间依赖算法，消除了子阵之间的依赖关系，实现了f列的并行消去和g列的并行消去。以9乘9矩阵为例，具体算法的过程如下：

1.将矩阵分块，如现有技术中的图1所示。

2.依次消去每块左下方元素(a2，a5，a8，a3，a6，a9)和上方元素(c1，c4，c7)，块内串行，块间并行，如现有技术中的图2所示。

3.消去每块上方元素(c3，c6)，存在块间数据依赖，块间串行，如图7所示。

4.依次消去每块左侧最下元素(f6，f9，g6，g3)，进行依赖数据预处理，块间串行。如图8所示。

5.顺序消去每块左侧元素和右侧元素(f5，f8，f4，f7，g1，g4，g7，c2，c5，c8)，块内串行，块间并行，如图9所示。

6.求解完毕，如图10所示。

上述的矩阵块间依赖消除算法，将原分裂法中具有数据块间依赖的部分进行预先消元，为后续块间并行消元打下基础，使得整个算法适用于GPU的并行处理

步骤602，将经过依赖消除的矩阵块间进行并行求解实现块间并行消元。

在本步骤中，经过依赖消除的矩阵分块之间可以并行处理，从而大幅提高计算性能的基础；但是在并行求解的过程中，需要考虑每个block中的线程数，每个线程负责的行数，使得单个线程所做工作尽量小。因此，针对为面向GPU加速求解三对角矩阵方程的特点，根据分块大小以及GPU线程数目，对将经过依赖消除的矩阵块间进行并行求解。

步骤603，利用主机内存和GPU显存的透明映射，使用主机CPU进行串行运算实现主机和GPU之间的数据传输通路。

在本步骤中，造成数据依赖的元素串行消去后可以实现大规模并行处理，但是串行的消元会造成大量的延时，为了进一步提高处理速度，采用CPU进行串行处理，利用一次串行操作消去造成数据依赖的元素而为了减少主机与GPU的数据传输时间。

采用映射存储器方式实现主机内存和GPU显存的透明映射，具体如下：

intcuda_device；

cudaDevicePropdeviceProp；

cudaGetDevice(&cuda_device)；

cudaGetDeviceProperties(&deviceProp,cuda_device)；

cudaSetDeviceFlags(cudaDeviceMapHost)；

cudaMallocHost((void**)&a,(N*sizeof(float)))；

cudaMallocHost((void**)&c,(N*sizeof(float)))；

利用主机内存和GPU显存的透明映射，一方面减少了数据传输量，另一方面简化了数据传输步骤，进一步挖掘了算法的整体性能。

本发明在实验平台上完成了验证，并达到了良好的实验效果。实验平台的具体配置是NVIDIAGT540m显卡(计算能力2.1，2SMs，2G显存)，2.30GHzInteli5-2410CPU(4核)，4GBDDR3RAM，操作系统使用CentOS6.0，内核版本2.6.32，CUDAToolkit版本4.0。

实验中与CPU的串行追赶法进行性能的比较，可以看出在大规模数据的情况下，加速效果还是比较明显的，实验结果下表1所示。

表1

问题规模	追赶法CPU串行	分裂法GPU并行	加速比
				1048576	482ms	54ms	8.93
26214400	780ms	12146ms	15.57

本发明中，面向GPU计算特点，充分挖掘矩阵数据分块之间的并行性，预先消除了矩阵分块间的数据依赖，将简单重复的块内消元置于GPU内部，利用大量的GPU核心进行并行处理，大大加快了三对角矩阵的求解速度。同时，利用主存与显存之间的映射机制，进一步提高的系统的运算效率。这种面向GPU并行的三对角矩阵方程求解算法所具有的上述优点，与传统的分裂法求解三对角矩阵的算法相比，本发明在保证了算法正确性的同时，大大降低了求解三对角方程的时间，以及本算法普适于各种GPU众核处理器架构。故本发明在实际数值分析实践中具有很高的技术价值和实用价值。

虽然本发明所揭露的实施方式如上，但所述的内容仅为便于理解本发明而采用的实施方式，并非用以限定本发明。任何本发明所属领域内的技术人员，在不脱离本发明所揭露的精神和范围的前提下，可以在实施的形式及细节上进行任何的修改与变化，但本发明的专利保护范围，仍须以所附的权利要求书所界定的范围为准。

Claims (5)

1.一种面向图形处理器GPU并行的三对角矩阵方程求解方法，其特征在于，包括：

将面向GPU并行的三对角矩阵进行矩阵块间依赖消除实现块内消元；

将经过依赖消除的矩阵块间进行并行求解实现块间并行消元；

利用主机内存和GPU显存的透明映射，进行主机和GPU之间的数据传输通路。

2.根据权利要求1所述的面向GPU并行的三对角矩阵方程求解方法，其特征在于，所述进行矩阵块间依赖消除，具体为：

根据矩阵数据分块之间的并行性，消除矩阵分块间的数据依赖，并将具有数据块间依赖的部分进行块内消元。

3.根据权利要求2所述的面向GPU并行的三对角矩阵方程求解方法，其特征在于，所述进行矩阵块间依赖消除，具体为：

将矩阵分块；

依次消去每块左下方元素和上方元素，块内串行，块间并行；

消去每块上方元素，存在块间数据依赖，块间串行；

依次消去每块左侧最下元素，进行依赖数据预处理，块间串行；

顺序消去每块左侧元素和右侧元素，块内串行，块间并行。

4.根据权利要求3所述的面向图形处理器GPU并行的三对角矩阵方程求解方法，其特征在于，所述进行并行求解，具体为：

根据分块大小以及GPU线程数目，对将经过依赖消除的矩阵块间进行并行求解。

5.根据权利要求4所述的面向图形处理器GPU并行的三对角矩阵方程求解方法，其特征在于，所述进行主机和GPU之间的数据传输通路，具体为：

利用主机内存和GPU显存的透明映射，使用主机CPU进行串行运算实现主机和GPU之间的数据传输通路。