基于CUDA的并行PCG方法自动生成系统研究与实现

1、硕士学位论文论文题目： 基于CUDA的并行PCG方法自动生成系统研究与实现作者姓名壬空 指导教师高家全教授 学科专业计算机技术 ， 培养类别 全日制专业学位硕士所在学院 计算机科学与技术学院浙江工业大学硕士学位论文基于CUDA的并行PCG方法自动生成系统研究与实现作者姓名：王宇指导教师：高家全教授浙江工业大学计算机科学与技术学院Dissertation Submitted to Zhejiang University of Technologyfor the Degree of MasterThe Automatic Generation System of the Parallel PCGM

2、ethod Based on CUDACandidate: Yu WangAdvisor: Prof. Jiaquan GaoCollege of Computer Science and TechnologyZhejiang University of TechnologyMar. 2017浙江工业大学学位论文原创性声明本人郑重声明：所提交的学位论文是本人在导师的指导下，独立进行研究工作 所取得的研究成果。除文中已经加以标注引用的内容外，本论文不包含其他个人或 集体已经发表或撰写过的研究成果，也不含为获得浙江工业大学或其它教育机构的 学位证书而使用过的材料。对本文的研究作出重要贡献的个人和集

3、体，均已在文中 以明确方式标明。本人承担本声明的法律责任。作者签名：日期:I产3月2g日学位论文版权使用授权书本学位论文作者完全了解学校有关保留、使用学位论文的规定，同意学校保留 并向国家有关部门或机构送交论文的复印件和电子版，允许论文被查阅和借阅。本 人授权浙江工业大学可以将本学位论文的全部或部分内容编入有关数据库进行检 索，可以采用影印、缩印或扫描等复制手段保存和汇编本学位论文。本学位论文属于1、保密口，在一年解密后适用本授权书。2、保密口，在二年解密后适用本授权书。3、保密口，在三年解密后适用本授权书。4、不保密SZ（请在以上相应方框内打“ J”）作者签名: 王亨日期：2W笋3月2了日导

4、师签名：而氟*日期：年J月左日基于CUDA的并行PCG方法自动生成系统研究与实现摘 要预条件共辘梯度(PCG)算法作为求解稀疏线性方程组的主流方法之一，近年来随着 问题规模的增大和GPU计算能力的快速提高，用于求解大规模问题的并行PCG算法引起 了更广泛的关注。并行PCG算法的研究重点是针对该方法的主要成分研究实现高效的并行 方法。如果通过人为方式寻找各个成分的最优实现，由于每个成分实现方式多样，而且影 响其性能的参数取值范围广，显然需要极大的工作量。为此，本论文通过优化建模技术，对PCG算法的主要成分分别构建并行优化性能模型, 从现有的核函数中选择出最优的核函数和参数配置，达到快速生成高效并

5、行PCG算法的目 的。本论文的主要工作和贡献如下：提出矢量运算、矢量内积的并行优化性能模型。分别对矢量运算、矢量内积建立并 行优化性能模型，通过决策树生成算法，自动生成决策树。实验证明，本论文的矢量运算、 矢量内积决策树对核函数以及参数的选择非常有效。提出稀疏矩阵矢量乘(SpMV)的并行优化性能模型。以5种经典的存储格式为例, 建立并行优化性能模型，并通过自动选择最优核函数算法，自动选择稀疏矩阵最佳的存储 格式、最优的核函数以及参数配置。实验证明，本论文的SpMV并行优化性能模型预测核 函数的执行时间的精度达95%以上，自动选择核函数算法具有鲁棒性，可靠性强。提出PCG并行优化框架。该框架包含

6、PCG算法各个成分的并行优化性能模型，各 个模型相互独立，具有很强的可扩展性。设计并实现基于CUDA的并行PCG方法自动生成系统。该系统使用图形化界面为 PCG算法的主要成分构建并行优化性能模型，自动生成高效的并行PCG方法。实验证明， 自动生成系统是有效可行的，通过本系统自动生成的并行PCG方法在单个GPU上的平均 加速比为56.91,在单节点2个GPU上的平均加速比可达到104.06o关键词：并行PCG方法，性能模型，自动生成系统，CUDA, GPUThe Automatic Generation System of the Parallel PCGMethod Based on CUDA

7、ABSTRACTThe preconditioned conjugate gradient (PCG) algorithm is one of the popular methods for solving large sparse linear systems. In recent years, accelerating the PCG algorithm on GPU has attracted considerable attention. However, on a specific multi-GPU platform, producing a highly parallel PCG

8、 implementation for any laige-sized problem requires significant time because several manual steps are involved in adjusting the related parameters and selecting an appropriate storage format for the matrix block that is assigned to each GPU.Therefore, using the optimizing model technology, we const

9、ruct the performance model for each one of main components of the PCG algorithm, and thus rapidly generate the parallel PCG algorithm by automatically selecting the optimal kernel and corresponding parameters from existing kernels. The main work and contributions are summarized as follows:Construct

10、the parallel optimization performance models for the vector operation and inner product. Utilizing the vector-operation and inner-product optimization models, decision trees are automatically generated.Construct parallel optimization performance model for SpMV. We take five classical storage formats

11、 and corresponding kernels to construct the performance models. Experimental results show that the accuracy of the execution time that is estimated by our proposed SpMV optimization performance model is more than 95%.Design a parallel optimization framework of PCG. In our proposed PCG optimization f

12、ramework, each model is independent and easily extensible.Implement an automatic generation system of the PCG method. This system can use the graphical visualization interface to build the parallel optimization performance model for each one of main components of the PCG algorithm, and thus automati

13、cally generate the PCG algorithm with high performance. Experimental results show that the average speedup ratios of the parallel PCG algorithm are 56.91 and 104.06 on one GPU and two GPUs, respectively.Key Words: parallel PCG method, performance model, automatic generation system,CUDA, GPU摘 要第1章绪 论

14、1.1课题研究的背景和意义 TOC o 1-5 h z HYPERLINK l bookmark64 o Current Document 国内外研究现状及趋势21.2.1并行SpMV算法的研究2并行PCG算法的研究3SpMV性能评估的研究31.3研究内容1.4论文章节安排相关技术及PCG并行优化框架 TOC o 1-5 h z HYPERLINK l bookmark91 o Current Document CUDA 介绍7CUDA并行计算7CUDA编程模型8CUDA存储器模型9 HYPERLINK l bookmark101 o Current Document 2.2稀疏矩阵存储格式1

15、0COO存储格式10CSR存储格式10DIA存储格式11ELL存储格式11HYB存储格式11 HYPERLINK l bookmark112 o Current Document PCG 算法12 HYPERLINK l bookmark125 o Current Document PCG并行优化框架14 HYPERLINK l bookmark128 o Current Document 本章小结15 HYPERLINK l bookmark133 o Current Document 第3章矢量运算和矢量内积并行优化性能模型研究16 HYPERLINK l bookmark136 o Cu

16、rrent Document 3.1矢量运算的并行优化性能模型16获取GPU特性163.1.2核模型16实验设置173.1.4并行优化性能模型的构建173.1.5生成决策树18 HYPERLINK l bookmark150 o Current Document 矢量内积的并行优化性能模型19获取GPU特性20核模型20实验设置203.2.4并行优化性能模型的构建203.2.5生成决策树22 HYPERLINK l bookmark161 o Current Document 本章小结22 HYPERLINK l bookmark166 o Current Document 第4章SpMV和预

17、条件子的并行优化性能模型研究24 HYPERLINK l bookmark169 o Current Document SpMV的并行优化性能模型24获取GPU特性24核模型25实验设置254.1.4并行优化性能模型建立274.1.5自动选择最优核函数算法32 HYPERLINK l bookmark208 o Current Document 4.2预条件子并行算法以及并行优化性能模型33 HYPERLINK l bookmark211 o Current Document 本章小结34 HYPERLINK l bookmark216 o Current Document 第5章系统实现与实

18、验比较35 HYPERLINK l bookmark219 o Current Document 5.1系统设计与图形化交互建模35系统设计35图形化交互建模37 HYPERLINK l bookmark236 o Current Document 实验比较435.2.1测试矢量运算、矢量内积决策树的有效性435.2.2测试SpMV并行优化性能模型预测核函数执行时间的精准度455.2.3测试自动选择最优核函数算法选择最优核函数的精准度48自动生成的并行PCG方法性能测试49 HYPERLINK l bookmark239 o Current Document 本章小结51 HYPERLINK

19、l bookmark244 o Current Document 第6章结论与展望53 HYPERLINK l bookmark247 o Current Document 结论53 HYPERLINK l bookmark255 o Current Document 展望54 HYPERLINK l bookmark261 o Current Document 参考文献55 HYPERLINK l bookmark322 o Current Document 致 谢59 HYPERLINK l bookmark325 o Current Document 攻读学位期间参加的科研项目和成果60

20、第1章绪 论1.1课题研究的背景和意义大型稀疏线性方程组的求解一直是科学与工程计算领域里最重要的问题之一。目前， 求解稀疏线性方程组的方法主要分为直接法和迭代法两大类。直接法是经过有限步四则 运算求得近似解，该方法对于低阶稠密矩阵方程组的求解比较有效。迭代法是经过迭代计 算，在迭代的过程中逐渐逼近于精确解。相比直接法，迭代法的运算量和存储量小而备 受研究者们的青睐。预条件共貌梯度（PCG）算法是迭代法中最流行的方法之一，近年 来，随着人们探索问题规模和深度的增大，迫切需要提高其计算效率。随着互联网的发展，大数据时代已经悄然而至。大数据带来的其中一个挑战就是要求 处理速度快囹，在海量的数据面前，

21、处理数据的效率就显得尤为重要。近年来，GPU已经 演变成一个具有高度并行、多线程的多核处理器，并且拥有强大的计算能力和极高的内存 带宽。2006年11月，NVIDIA公司发布了通用并行计算平台和编程模型CUDA，它允许 程序员使用高级语言进行CUDA编程，并且还提供了大量并行库，例如用于基本线性代数 的cuBLAS和用于加速训练深度神经网络的cuDNN等，使GPU编程变得非常容易。由于GPU的高计算能力，研究者为了提高PCG算法的计算效率，将PCG算法迁移到 GPU平台上。目前，GPU加速的并行PCG算法已有较多的研究成果，例如并行MIC预条 件PCG算法A和并行SSOR预条件PCG算法等。不

22、难发现，PCG算法主要由矢量运 算、矢量内积、预条件子方程求解和稀疏矩阵矢量乘（SpMV）等主要成分组成。如果想 获得高性能的并行PCG算法，那就必须研究这些主要成分的高效并行实现方法。高效的CUDA程序不仅要求程序要充分使用GPU稀有的片上资源（寄存器、共享内 存等），使其利用率达到最大化，而且在运行核函数时，还需要人工进行合理地分配线程, 使得程序并行达到最大化。不同的GPU,它达到最佳性能的参数往往是不同的，需要遍历 块内线程数，以找到最合适的线程分配。对于PCG算法来说，现有GPU加速的矢量运算、 矢量内积、预条件子方程求解和稀疏矩阵矢量乘（SpMV）方法较多，特别是SpMV,还 与存

23、储格式相关。因此，如果生成并行PCG算法的过程中使用人工参与方式进行调优和选 择最优的核函数，当面对计算机集群时，就会变得繁重，需要极大的工作量。这促使了我 们的动机：通过优化建模技术，对PCG算法的主要成分分别构建并行优化性能模型，以达 到能自动从现有的核函数中为PCG算法的主要成分快速选择出最优的核函数和参数配置, 自动生成高效的并行PCG算法，避免大量繁琐的人工调优过程。因此，本论文研究不是构造一个新的核或算法，而是从PCG算法的主要成分出发，通 过构造其与问题无关的优化模型，达到自动快速生成高效并行PCG算法的目的，这与其他 的PCG算法加速研究有着本质的区别。1.2国内外研究现状及趋

24、势与并行PCG方法自动生成系统相关的研究包括并行SpMV算法、并行PCG算法和 SpMV性能评估等方面，下面从这几个方面分别介绍其国内外研究现状。1.2.1并行SpMV算法的研究如果一个矩阵中的零元素个数远远多于它的非零元素个数，那么这样的矩阵称之为稀 疏矩阵。对于稀疏矩阵，为了减少不必要的冗余计算和存储空间，通常需要对稀疏矩阵进 行存储。目前，存在多种稀疏存储格式，每种存储格式都有自身的特点，对同一个稀疏矩 阵来说，采用不同的存储格式就会获得不同的SpMV性能。对于稀疏矩阵而言，只有根据 它的非零元素特征分布，选择最合适的存储格式，SpMV核函数才能发挥最佳性能。由于在迭代方法中，SpMV占

25、有重要的地位I%研究者们着重对它进行了研究。2008 年，Bell和Garlandtl3详细分析了 COO、CSR、ELL和DIA这4种经典稀疏存储格式的优 缺点，并提出HYB存储格式，将ELL和COO存储格式一起使用，减少了 ELL存储格式 零元素的填充。基于这些存储格式，作者们提出了一些高效的SpMV CUDA核函数，并把 它们封装在CUSP开源包中供研究者们下载使用。后来研究者们又对这些存储格式进行了改进和扩展。Zheng等Ml对ELL进行了改进， 提出了 BiELL存储格式，可以减少零元素的填充，并在GPU上基于BiELL存储格式实现 了 SpMV,当矩阵每行非零元素不均匀的时候，Bi

26、ELL存储格式性能要比ELL好o Maggioni 等I*在ELL基础上提出了一种自适应的ELL存储格式：AdELL,使得每个warp计算时负 载均衡。Dang等提出sliced COO (SCOO),进而实现一种SCOO SpMV核，当矩阵是 单精度浮点数时，SCOO SpMV性能要比COO和CSR SpMV核都要好。Liu等刀为了克 服CSR线程负载不平衡的缺点，设计了一种CSR5存储格式，提出了高通量的SpMV核。 为了减少存储空间，Yan等18使用分块技术，提出了 BCCOO和BCCOO+存储格式，有效 减少了传统COO的存储空间。Choi等19提出了 BCSR存储格式，减少了 CSR

27、行索引和列浙江工业大学硕士学位论文 索引的存储空间，利用基于模型的自动选择参数框架，自动调节块block大小。当选择合 适块大小(Blocksize)的时候，BCSR SpMV核的性能比CSR SpMV核优。Tang等刖使 用位表示优化技术压缩index和data数组来减少存储空间，提出BRO-ELL、BRO-COO、 BRO-HYB 等 SpMV 核。另外，研究者们还提出很多新的存储格式，例如CSX存储格式21, BRC存储格式22, BRO存储格式，AMB存储格式网，SELL-C-o存储格式和JAD存储格式网等，基于 这些存储格式，也提出了一些有效的SpMV核。1.2.2并行PCG算法的研

28、究Jacobi预条件子是较早提出用于CG算法的预条件子2728,为了提高其效率，研究者 又设计了 ILU预条件子29、SSOR预条件子3、代数多级网格预条件子8刖和Gauss-Seidel 预条件子皿等，进而提出了许多有效的PCG算法。由于PCG算法中预条件子方程求解时, 需要前推和回代过程，导致不易并行化，为提高其效率，有研究者对其进行了研究。Liu和Chen等闵钏在GPU上实现了多个预条件PCG并行算法，实验数据显示可以有 效求解大型线性系统。Li和Saad。】实现了 MIC预条件PCG并行算法，获得了 3倍的加速 比。Gao和Liang等in对稀疏对称正定的七对角矩阵，提出了一种有效的M

29、IC预条件GPU 并行实现方法，比cuSPARSE库实现快3倍左右。为了避免预条件子方程求解过程中的前 推和回代，也有研究者利用稀疏近似逆技术，将预条件子方程求解转变成仅需一个SpMV 的运算，提高并行效率。例如，Helfenstein和 险村提出了 SSSOR近似逆预条件PCG算 法，获得了 10倍的加速比。王志超等D习利用诺依曼多项式分解技术，给出了在GPU上一 阶和二阶SSOR稀疏近似逆并行PCG算法。1.2.3 SpMV性能评估的研究性能评估是本论文模型建立的核心，直接影响模型预测的精准度。Monteiro等使用 矩阵样本集合训练静态模型(STOMP),通过静态模型预测SpMV的执行时

30、间，该静态模 型预测精度达到95.3%。Neelima等闵闵提出稀疏矩阵存储格式预测模型，通过分析稀疏矩 阵非零元素的分布特点，按照事先制定好的规则进行预测，如果有不满足规则的稀疏矩阵, 需要结合CPU跟GPU的传输开销再进行判断，但该模型在预测稀疏矩阵存储格式的过程 中并未考虑到SpMV计算的时间。Li等39,4。利用概率模型，可以准确描述稀疏矩阵的非零 元素分布情况，能预测稀疏矩阵存储格式的选择，但每次预测前都需要对目标稀疏矩阵进 行概率模型分析，再根据GPU硬件参数进行性能评估。Zardoshti等叩】提出一种自适应的 -3-运行时系统来选择稀疏矩阵最佳的存储格式，该系统需要将目标矩阵划

31、分成几个样例矩阵 输入系统中，系统根据样例矩阵来测试和对比，最后输出最佳存储格式。GUO等42,43通过 构造基矩阵集合，提出了一种基于profile的GPU SpMV性能模型，能预测CSR、ELL、 COO和HYB SpMV核函数的执行时间，该模型只跟GPU资源本身有关，它只需要一次建 模，后面不管有多少测试矩阵，都不要再重新建模，但论文中没有考虑到GPU参数（线 程分布等）对SpMV性能的影响。Guo和Lee】44】对Guo等的模型做了进一步改进，增加 了 Li等I的矩阵分析方法，但准确率还是跟原来一样。本论文是受到Guo等性可性能模型的启发，在此基础上进行优化、扩充，并作进一步更 深的研

32、究，提出了只跟GPU资源有关并且具有很强扩展能力的并行优化框架。1.3研究内容本论文研究目标并不是为了提出一种新的PCG并行加速算法，而是系统根据所需要的 运算及给定的问题，通过并行优化建模技术，构造矢量运算、矢量内积、稀疏矩阵矢量乘 （SpMV）和预条件子方程求解等优化模型，从现有的核函数中挑选出最优的核函数，自 动生成高效的PCG并行算法。本论文的核心是结合CUDA特性，分析影响PCG算法性能的关键成分：矢量运算、 矢量内积、稀疏矩阵矢量乘（SpMV）和预条件子方程求解，分别对这些关键成分构建并 行优化性能模型。构建的模型只与GPU本身资源有关，对一种类型的GPU,我们的模型 仅需构建一次

33、。本论文的自动生成系统具有很强的可扩展性，假设没有包括在系统框架中 的核模型，只要能建立起性能模型，就能将它加入到系统框架中。下面是本论文的具体研 究内容：矢量运算的并行优化性能模型研究对于矢量运算而言，影响它们的性能好坏主要是核函数的选择以及线程资源分配。通 过建立矢量运算的并行优化性能模型，最后生成决策树。给定一个矢量，通过决策树能够 自动快速地找出最优的核函数以及运行时的线程参数分配，使得矢量运算核函数性能达到 最优。矢量内积的并行优化性能模型研究由于矢量内积运算存在归约操作，我们定义最小线程块数来权衡线程块数大小。构建 矢量内积的并行优化性能模型，最后生成决策树。给定一个矢量，通过决策

34、树快速找出线 程数配置，使得矢量内积核函数性能达到最优。SpMV的并行优化性能模型研究SpMV的情况要比矢量运算和矢量内积复杂。除了要合理分配线程外，还要找出最合 适的存储格式，只有选择最合适的存储格式，SpMV核函数才能发挥最佳性能。本论文以 5个经典的稀疏存储格式CSR、DIA、ELL、COO和HYB以及它们对应核函数为例，建立 SpMV的并行优化性能模型，通过自动选择最优核函数算法，给定一个稀疏矩阵，能够快 速找出最适合的存储格式、核函数以及线程分配。预条件子的并行优化模型研究通过建立的近似逆预条件子，预条件子方程求解转变成为SpMV,这样在预条件子求 解过程即可使用SpMV并行优化性能

35、模型来调优程序性能。PCG的并行优化框架提出PCG并行优化框架，该框架包含各个运算的并行优化性能模型，具有很强的可扩 展性。并行PCG方法自动生成系统设计与实现并行PCG方法自动生成系统，使用图形化界面操作，为PCG算法的主要 成分构建并行优化性能模型，选择最优核函数，自动生成高效的并行PCG算法。1.4论文章节安排第一章绪论首先介绍基于CUDA的并行PCG方法自动生成系统研究与实现的研究背景，以及 研究的目的和意义。其次介绍并行SpMV算法、并行PCG算法和SpMV性能评估这三方面的国内外研 究现状及趋势。介绍本文的主要研究内容。第二章相关技术及PCG并行优化框架介绍本论文的相关技术，CUD

36、A、5种经典的稀疏存储格式和PCG算法。提出PCG并行优化框架。第三章矢量运算和矢量内积的并行优化性能模型研究详细介绍矢量运算、矢量内积的并行优化性能模型的构建过程，以及决策树生成算法。第四章SpMV和预条件子的并行优化性能模型研究SpMV的并行优化性能模型研究按照PCG并行优化框架，详细介绍了获取GPU特性、核模型、实验设置、各个存储 格式的并行优化性能模型建立以及自动选择最优核函数算法这五个模块。预条件子的并行优化性能模型研究使用SSOR稀疏近似逆预条件子，将预条件子方程求解过程变成SpMV运算。第五章系统实现与实验比较详细介绍系统设计，通过ELL存储格式案例来介绍系统图形化建模的过程。实

37、验比较，分别从以下四个方面进行实验对比：1）测试矢量运算、矢量内积的决策树有效性；2）测试SpMV并行优化性能模型预测 核函数执行时间的精准度以及跟其他研究者模型对比；3）测试自动选择最优核函数算法 选择最优核函数的精准度；4）自动生成的PCG方法性能测试。第六章结论与展望。对本文的研究工作进行总结，并展望未来的研究内容。第2章 相关技术及PCG并行优化框架本章主要有两个方面的内容，一、介绍本论文涉及的相关技术，主要包括CUDA、稀 疏矩阵存储格式和PCG算法。二、提出本文的PCG并行优化框架。CUDA 介绍CUDA并行计算随着计算机芯片制作工艺的不断提升，现代计算机已经演变为多核、多线程处理

38、器， 这就要求编程人员要从以前的串行、单线程的问题求解方法切换到多线程并行执行的问题 求解方法4习。并行计算的产生得益于多核多线程处理器的出现，它通常要把一个大问题分 解为很多独立的小问题，使用多个处理器或者计算机同时进行计算I。并行计算最主要的 目的就是提高计算的速度。近年来，GPU凭借着高计算强度和高内存带宽在并行计算领域中如火如荼。在某种程 度上讲，也得益于成熟的编程模型。2006年11月，NVIDA提供了一个易用的编程接口， 也就是通用并行计算平台和编程模型CUDA。CUDA是在标准语言的基础上进行了扩展， 使得程序员只要学习目前流行的编程语言（C语言等），就能够非常容易地使用它。为了

39、让 不同领域的人都能简单地使用CUDA技术，NVIDA公司还发布了很多已经封装好的并行 库，如:cuBLAS, cuSPARSE48等。DRAMDRAMGPUCPU图2-1 CPU结构与GPU结构GPU与CPU结构有很大区别，如图2-1所示。从图2-1中，可以非常直观地发现GPU 的处理器（绿色部分）要比CPU多，而CPU的控制器（黄色部分）和高速缓存（红色部 分）要比GPU多。可见，GPU设计了更多的处理器去处理数据而不是数据缓存和逻辑控 制，是一个面向计算、吞吐量的多线程、多核处理器。CUDA编程模型在GPU上实现程序并行化一般需要满足以下三个条件：a）细粒度并行；b）计算密集 型；c）各

40、计算任务之间松耦合，最好相互独立。CUDA编程模型允许程序在异构系统上执 行，通常把串行代码放在主机端（host）执行，并行代码放在设备端（device）执行。把在GPU 设备上的运行代码称为核函数（kernel），使用_global_关键字声明。核函数一旦被调用, 控制权马上会返回给主机端（host）, GPU和CPU交替执行。如图2-246所示。CUDA C ApplilcationHost=CPUHost codeParallel kernelDevice=GPUHost codeHost=CPUParallel kernelDevice=GPU图2-2 CUDA CPU与GPU交替执行

41、CPU调用核函数时，必须按照以下语法：Kernel_function （parml, parm2, .）其中，gridDim表示网格内线程块数量和维度，blockDim表示线程块内线程数量和维 度，Ns用于动态分配共享内存时指定的空间大小，s指定调用核函数时对应的流。CUDA内核以线程块构成的网格（grid）进行启动，一个网格由许多线程块（block） 构成，每个线程块（block）由许多线程（thread）组成，如图2-3所示。线程块内的每个线程共享同一个线程块索引(blockldx),它是CUDA内置的uint3类型的变量，作为线程 块在网格内的唯一索引。线程块内的每个线程也有内置的uin

42、t3类型变量threadldx,作为 线程在线程块中的唯一索引，通过访问thieadldx和blockldx很容易区分每个线程。CUDA 中网格和线程块都是三维的，通过dim3类型的CUDA内置变量gridDim和blockDim来指 定。1111 /七 ，/1111、 ;、1 、1gridDim.x（网格） AEapC8blockDim.x(线程块)图2-3 CUDA线程层次结构CUDA把32个线程称为一个waip (线程束)，并作为它的基本执行单元。一个warp 里的线程执行同一条指令，因此需要减少或者避免程序中判断语句的出现，比如if, for 语句等，防止出现warp分支。因为一个wa

43、rp里的不同线程执行不同的路径，会导致程序 性能下降。CUDA存储器模型CUDA存储器模型提供了很多可编程的内存类型：寄存器(Registers)本地内存(Local memory) 共享内存(Shared memory) 全局内存(Global memory) 常量内存(Constant memory)和质地内存(Texture memory),编程人员可以自由地控制数据的存储。每个内存 类型都有不同的使用范围、生命周期以及缓存行为，如表2-1所示。寄存器是GPU稀有的 片上资源，是GPU最快的存储器。寄存器变量是每个线程私有的，当线程执行结束，寄 存器变量就会失效。共享内存也是GPU片上资

44、源，它跟CPU的L1高速缓存相似，但共 享内存是可编程存储器，可以被同一个块内的线程访问，在使用的过程中需要避免bank 冲突的产生，防止程序性能下降。在编写CUDA程序时，需要充分利用GPU的片上资源, 才能获得更高的加速比。表2-1各种存储器比较存储器位置拥有缓存生命周期寄存器GPU片上N/A与Thread相同本地内存板载显存无与Thread相同共享内存GPU片上N/A与bock相同常量内存板载显存有主机程序配置质地内存板载显存有主机程序配置全局内存板载显存无主机程序配置2.2稀疏矩阵存储格式由于稀疏矩阵存在大量的零元素，为了节省存储空间和减少计算冗余量，需要对矩阵 做压缩处理。近年来，学

45、者们对稀疏矩阵存储格式进行了很多的研究，也做了很多改进。 常见经典的稀疏矩阵压缩格式有：COO、CSR、DIA、ELL和HYB等。COO存储格式COO (Coordinate Format)是我们最简单、通用的一种存储格式网。这种存储结构使 用row, col, data这三个数组把稀疏矩阵的每个非零元素值保存下来，其中mvv数组存储非 零元素对应原始矩阵的行索引，c。/数组存储非零元素对应原始矩阵的列索引，也幻数组 存储对应的非零元素值。这三个数组的长度由稀疏矩阵总的非零元素个数决定的。COO存 储格式具有通用性、简单灵活易于操作，它可以存储任何类型的稀疏矩阵，但我们需要存 储每个非零元素的

46、坐标和非零值，需要较大的存储空间。具体例子如图2-4 (b)所示。CSR存储格式CSR( Compressed Sparse Row Format)在COO基础上进行了改进，使用行偏移数组ptr 来代替COO的mw数组，列数组必力ces和非零元素值数组出亿存储方式跟COO 一样。行 偏移数组P”保存也如数组每行第一个元素的起始偏移位置。假设稀疏矩阵有N行，则ptr 数组长度为N+1。稀疏矩阵第，行的非零元素个数可以通过计算p tri+l-ptri来获得。 CSR同样具有COO的灵活、易操作的特点，同时要比COO存储空间少，是存储格式中比 较常用的一种。具体例子请参考如图2-4 (c)。DIA存

47、储格式DIA (Diagonal),对角线压缩存储法，按对角线形式存储，它使用次?柩和必M数组 来表示，其中，也幻数组存储的是矩阵的非零元素的值，也2,数组存储的是子对角线相 对于主对角线的位移。当咖e心0时，则表示该非零元素位于主对角线上方，距离为 offseti的子对角线上。当offseti0时，则表示该非零元素位于主对角线下方，距离为 offsets的子对角线上。当offseti = 0时，表示该非零元素在主对角线上。DIA存储格 式适合对角性很好的稀疏矩阵，执行SPMV时，要比CSR、COO效率高。由于DIA存储 格式的特殊特点，导致它不具有通用性。具体例子请参考如图2-4 (d)。E

48、LL存储格式ELL (ELLPACK),它由z所此攻和出幻这两个数组来表示。对于一个MxN的稀疏矩 阵，我们把稀疏矩阵每行非零元素个数的最大值记为K。ELL存储格式分别用一个MxK 的也幻数组和z如7诂绐数组来存储非零元素值和对应的列索引，那些每行平均非零元素个 数比K小的行用零元素填充。当稀疏矩阵每行的非零元素个数较均匀的时候，就推荐使 用ELL存储格式。如果每行平均非零元素个数跟K相差较大，则会占用不必要的存储空 间和冗余计算，导致性能下降5。具体例子请参考图2-4 (e)。HYB存储格式由于ELL存储格式在存储每行非零元素个数分布不均匀的矩阵时，会降低程序性能。 Belltl3结合ELL

49、和COO的特点，提出了一种HYB的存储格式。通过计算H值，把稀疏 矩阵每行H个非零元素采用ELL存储格式来存储，超出H部分的非零元素则采用COO存 储格式存储，其中H的取值标准是：矩阵包含H个非零元素个数的行至少要占稀疏矩阵总 行数的三分之一I%由于ELL执行效率要比COO高，所以H值的选择对HYB来说是至 关重要的。我们应该把稀疏矩阵的主要部分使用ELL存储，剩余少量部分使用COO来存 储。具体例子请参考图2-4 (f)。这里对稀疏存储格式进行简单的总结：1、DIA和ELL存储格式在执行SpMV操作时效率最高。2、COO和CSR存储格式比ELL和DIA灵活，易于操作，通用性强。3、HYB存储

50、格式将ELL和COO的特点进行了结合，但由于ELL的执行效率要比 COO高的多，因此，应该合理选择H值，使得矩阵的绝大部分采用ELL存储格式，剩余 小部分采用COO存储格式。3070690208400051(a)稀疏矩阵row = 03ptr = 0 29col = 03indices = 03data = 32 8 4 5 1data = 3 71-*37-_02*-_37*-692offset = -1 0 2013692data =84*indices =12*data =84*51*23*51*(b) COO存储格式(c) CSR存储格式(d) DIA存储格式(e) ELL存储格式-3

51、7-0269indices =018412_51_23_data =ELL部分row = 1 co/ = 3data = 2(f)COO部分HYB存储格式图2-4稀疏矩阵表示2.3 PCG算法共轴梯度法(CG)是介于最速下降与Newton法之间的迭代方法，为了进一步提高收敛的速度，研究者们引入预处理矩阵，大大提高了它的迭代收敛速度5此 比如：ICCG算 法52-54、ILUCG算法26和sSORCG算法52,55等。本论文使用的是Chronopoulous等56提 出来的PCG算法，因为它把矢量运算和内积运算聚集在一块，这样可以有效减少核函数的 数量，具体算法步骤见算法2-1。算法2-1： P

52、CG算法Input: A, b, , x (initialize x with zeros)Output: xr b Ax co = Mxr s = AcoPo = Fco； /z = sTa); a = Pq3, /? = 0for it 1, 2,., MAX_ ITER dop = + = p; q = s + gqx x+ap, r raqif |r|7 thenbreak;endco = Mxr, s Acopx rT/ = sTcoIL 6 = PjPo； a=pJ(N-pBla)； Po=P12. end无论什么版本的PCG算法，它的运算主要包括SpMV、预条件子方程求解、矢量内

53、积 运算和矢量运算等主要成分。假设稀疏矩阵A(x)的每行平均非零元素个数为k, PCG 算法运行迭代次数为m ,则每个运算的计算复杂度如表2-2所示。表2-2各个运算的计算复杂度操作复杂度计算次数SpMVo(2kn)m + 1预条件子求解o(2kn)m + 1矢量内积运算o(2n)2m+ 2矢量运算。()4m+ 1从表2-2可以看出，SpMV和预条件子方程组的求解占PCG算法的绝大部分时间。如 果k/?)次核函数，对应的执行时间为：7；和号，定义核函数的平均执 行时间为：丁二亿一功/(a一”)。3.1.4并行优化性能模型的构建给定块内线程数小、任意矢量大小，计算出ne = nntx Dx,泌=

54、/(ex小)。 当块内线程数小一定时，同一个测试域里的矢量集对应的再值总是相同的，这意味着在 同一个测试域中它们使用的是同一个核函数。通过实验发现，在一个测试域中，当块内线程数小一定时，矢量大小跟核函数执行 时间存在一种线性关系，即：(3-1)其中，nr = 128, 256, . , Bf,，表示第，个测试域。例如：基于NVIDIA K40c,在测试域（128x3x65535, 128x4x65535上，取小的值为：128, 256, 512, 1024,得到矢量大小和核函数执行时间的线性关系曲线，如图3-1所示 （x轴表示矢量尺寸，缩小了 10倍；y轴表示核函数的执行时间）。xxXX图3-

55、1矢量运算关系曲线3.1.5生成决策树通过建立上面的并行优化性能模型，最终可以生成决策树。通过构建的决策树， 任给一个矢量尺寸，决策树都能够达到自动、快速、准确地找到最优核函数以及运行 时对应的参数配置。具体决策树生成算法如下所示。算法3-1矢量运算决策树生成算法遍历实验设置中的每个测试域（由,川（，=1,2.），执行以下操作:a）通过公式（3-1）拟合出每个测试域的线性直线，求出任何两条直线之间的 交点，并把交点保存在集合S中。b）如果5 = 0或者S集合中只有一个元素：为或者dj+i ,那么在测试 域中，最佳的块内线程数灯可以通过计算argmin”,7；岩 获得，其结果保存在 集合 P =

56、邑,dj+nt, nej 中。否则，把为和d州加入到S集合中，将S中的元素进行升序排序，对于每个 域(，必+1，(&e S , A = l,2,.,|s| , S S2 .七)，通过计算 argmin*；获 得最佳的块内线程数其结果保存在集合P = sk, s,+t, nt, ne中。输出集合P,即生成的决策树。例如，基于NVIDIA Tesla K40c,生成矢量运算的决策树如图3-2所示。图3-2矢量运算决策树3.2矢量内积的并行优化性能模型矢量内积运算跟矢量运算不同之处在于它包含归约操作。因此，我们单独对它进行了 研究，并提出一种基于profile的矢量内积并行优化性能模型，如图2-5的

57、线路b所示。它 主要包括：获取GPU特性、核模型、实验设置、并行优化性能模型的构建和生成决策树 这几个步骤。通过构建的决策树，给定任意大小的矢量，决策树能够自动、快速帮你配置 参数，使得程序性能达到最优。3.2.1获取GPU特性对于矢量内积运算，需要获取线程块内最大的线程数？，GPU全局内存最大字节数 G,网格X维度上最大的线程块数GPU流处理器个数Nsm,流处理器中32位寄存 器的最大个数Neg ,流处理器最大共享内存字节数Nmem ,流处理器最多线程块数N以及 流处理器允许最多线程数N等参数值。3.2.2核模型核模型里包含着不同运算的核函数。对于矢量内积，本论文使用的是Gao等in论文里

58、的矢量内积核函数。3.2.3实验设置实验设置主要的任务是找出所有可能的块内线程数力和构建benchmark矢量集。矢 量内积运算的实验设置跟矢量运算相似，这里只作简单介绍，具体可参考矢量运算的实验 设置。块内线程数小= 128, 256, 512, . , BL测试域设置为(1024,102400, (102400,1024x1024, (1024x1024,128x0,(128x0、,2x128x0,. ,(Bf xDAniax,其中，Nmx =Gm/(3x sizeof (double) o当混一定时，在一个测试域的ne = n/(ntxDx)总是相同的。如果对某个测试域取若干 个值作为矢

59、量尺寸，分别进行赋值操作，便形成一个benchmark矢量集。benchmark矢量中元素的值不会影响到矢量内积并行优化性能模型的性能，可以 根据均匀分布U0.5,1.5随机产生。3.2.4井行优化性能模型的构建由于内积核函数包含归约操作，核函数运行结束后会产生泌个局部归约结果(partial result),需要把泌个局部归约结果传到CPU主机端进行二次归约操作。由于数据通过PCI 插槽传输到CPU的代价是非常昂贵的，所以线程块数泌不能设置的太大。但是另一方面, 在CUDA程序中，线程块数渺值越大，意味着程序并行化能力越强。所以必须要对展的取值进行衡量，定义网格的最小线程块数油混为：NBnt

60、 = Nsm x min(Nreg/N；e8, Nmem/N,；,em, Ntd/nt, Nrb)(3-2)其中，N笋表示每个线程块需要的寄存器数量，NU表示每个线程块需要的共享内存 字节数。实验数据表明，当块内线程数小一定，取线程块数nb = NBlt时，核函数性能不一定是 最优的。为了使最小线程块数NBf的使用有更好的鲁棒性，定义线程块数 nb = NBntxi, 1 = 1,2,.,20,用户可以自己选择，的取值。通过实验发现：当位和泌一定时，在测试域中矢量大小跟核函数的执行时间存在着 一种线性关系：T：b,j = K,j ()(3-3)其中，nt = 128, 256, . , BT,