大规模稀疏线性方程组的GMRES—GPU快速求解算法
来源:测品娱乐
第23卷第4期 2011年4月 计算机辅助设计与图形学学报 Journal of Computer~Aided Design 8L Computer Graphics Vo1.23 NO.4 Apr.2011 大规模稀疏线性方程组的GMRES—GPU快速求解算法 柳有权 ,尹康学”,吴恩华 。’ l (长安大学信息工程学院西安710064) 北京100190) 2 (中国科学院软件研究所计算机科学国家重点实验室’(澳门大学科技学院澳门) (youquan@chd.edu.cn) 摘要:重开始广义极小残量法(GMREs)是求解大规模线性方程组的常用算法之一,具有收敛速度快、稳定性好等 优点.文中基于CUDA将GMRES算法在GPU上进行并行算法实现,尤其针对稀疏矩阵矢量乘法运算,通过合并访 问和共享内存策略相结合的手段使得算法效率大幅度提升.对于大规模数据集,在GeForce GTX 260上的运行结果 相对于Intel Core 2 Quad CPU Q94O0@2.66 GHz得到了平均4O余倍的加速效果,相对于Intel Core i7 CPU 920@ 2.67 GHz也可得到平均2O余倍的加速效果. 关键词:CUDA;GPGPU;重开始广义极小残量法;稀疏矩阵矢量乘法 中图法分类号:TP391 Fast GMRES-GPU Solver for Large Scale Sparse Linear Systems Liu Youquan , ,Yin Kangxue”,and Wu Enhua ’。 (School of Information Engineering,Chang an University,xi'an 710064) (State Key Laboratory of Computer Science,Institute of Software,Chinese Academy of Sciences,Beijing 100190) 。 (Faculty f oScience and Technology,University of Macau,Macao) Abstract:As a popular iterative method to solve linear equations,restarted generalized minimal residual method(GMRES)has the advantages of fast convergence and good stability.This paper implements a parallel GMRES in GPU based on CUDA.Particularly,the sparse matrix vector multiplication is optimized with coherence visiting and shared memory,which significantly improves the performance.We tested the paralleled GMRES on a GPU of GeForce GTX260,and compared its performance with those of the traditiona1 GMRES on Inte1 Core 2 Quad CPU Q9400@2.66GHz and Intel Core i7 CPU 920@2.67GHz,which showed 40 times of speed—up and 20 times of speed—up on average respectively. Key words:CUDA;GPGPU;generalized minimal residual method;sparse matrix vector multiplication 计算机图形学领域和一些实际工程应用中存在 很多偏微分方程的求解,如软体变形、流体仿真、几 的问题.这类求解通常采用迭代法进行数值计算,因 此此类方法的高效求解对这种复杂问题有着非常重 要的意义.该线性方程组可统一表示为Ax=b;其中 A为n× 大小的系数矩阵,X为 元变量,b为已知 何处理,这些方程在经过离散化后都转化成线性方 程组,从而将复杂问题求解变成一个可计算机求解 收稿日期:2010 09 25;修回日期:2010—11-30.基金项目:国家自然科学基金(60973066,60833007);中国科学院软件研究所计算机科学 国家重点实验室开放基金(SYSKF1004).柳有权(1976一),男,博士,副教授,CCF会员,主要研究方向为计算机图形学、虚拟现实.尹康学 (199o一),男,在校学生.吴恩华(1947),男,博士,教授,博士生导师,CCF高级会员,主要研究方向为计算机图形学、虚拟现实. 554 计算机辅助设计与图形学学报 第23卷 量.在目前诸多求解大规模线性方程组问题的迭代 法中,重开始广义极小残量法(generalized minimal residual method,GMRES)①_】 是很受欢迎的算法 之一,它通过Krylov子空间矢量的最小残量来迭代 求解,具有收敛速度快、稳定性好等优点. 目前有很多研究人员侧重于改进GMRES算 法,以进一步提高该方法迭代的效率.如全忠等口 通 过构造多项式预处理因子来克服GMRES算法有 时收敛很慢或停滞的缺陷,Habu等_4 通过调整重 新开始来加快GMRES的收敛速度. 但除了从算法层面来改进整体收敛速度,单步 计算的效率提升同样重要,这一方面依靠硬件的计 算能力的提升,另外新的计算架构通过对算法进行 并行化处理也可提升算法计算效率.传统的高性能 计算依赖大型机和计算集群,然而这样的计算系统 都很昂贵.GPU的发展为高性能计算提供了另外一 种思路 ],它采用众核架构,即芯片上集成了多个 并行处理单元.随着可编程性的出现,GPU从单纯的 图形流水线渲染转向通用计算上的应用(GPGPU), 现在在高性能计算机领域也开始崭露头角.GPU单 位计算成本的下降引起了很多研究机构和企业的 广泛重视,如国产天河系统由于采用CPU+GPU 以较低的代价获得非常高的性能,在最近的全球高 性能计算机排行榜上名列第一.NVIDIA[7 推出的 CUDA(compute unified device architecture)架构由 于编程方式的革新更是推动了GPU芯片在高性能 计算上的应用. 目前已有一些GMRES利用CUDA进行加速 的工作,如Wang等在GeForce GTX280图形卡上 获得20倍加速 ],Velamparambil等类似的工作在 GeForce 8800上获得13倍的加速LgJ,Ghaemian等 基于NVIDIA Tesla C870 GPU只获得60 的加 速口 .文献En]对作为各种迭代求解方法都要使用 的核心算法——稀疏矩阵与矢量乘法运算做了详细 分析.虽然上述工作各自的硬件平台略有不同,但整 体加速效率仍有提升的空间. 本文基于CUDA将GMRES算法在GPU上进 行重新设计,尤其是对稀疏矩阵与矢量乘法部分,通 过合并访问和共享内存的分配来优化负载;并充分 利用GPU的众核处理能力,使得算法效率相对于 CPU算法有大幅度提升.另外,通过跟NVIDIA提 供的稀疏矩阵与矢量相乘的GPU算法[1妇比较,本 文算法实现相对于NVIDIA自身的代码也有好几 倍的效率提升,同时代码在http://imlab.chd.edu. cn/cuda/ ̄公开,供研究人员免费使用. 1 GMRES算法分析 关于GMRES算法的详细描述请参考文献[2], 本文为了完整起见,对其稍作介绍.对于线性方程组 Ax=b,GMRES算法的m阶Krylov子空间为K埘一 span(b,Ab,A b,…,A一 6);GMRES通过求使残量 Ax 一6最小的矢量X ∈ 来逼近Ax—b的精确 解.但是,矢量b,Ab,A b,…,A一 6几乎是线性相 关的,因此通常采用Arnoldi迭代方法来找出正交 矢量', ,',2,…,v 作为m阶Krylov子空间的基.故 矢量 ∈K 可写成X 一V Y ,其中 ∈ 且y 是由',1,’,2,…,', 组成的 ×m矩阵. 通过Arnoldi迭代过程也可产生一个(m+1)×m 阶的上Hessenberg矩阵H 满足AV 一’,珥+1H .因 为V 是正交的,因此有ll Ax 一b lI—l lH Y 一 ll;其中e1一(1,0,0,…,o)是R啪+ 的标准基的 第一个矢量,并且 一ll Ax。一b Il,X。是初始矢量 (通常是零矢量).因此,求使得残量r =Ax 一b范 数最小的 ,就变成求r埘一H Y 一 范数最小的 问题,即为一个 阶线性最小二乘问题,通常情况 下m《 . 通过前面的描述可知,GMRES算法在迭代的 每一步中要进行如下操作: Step1.做一步Arnoldi迭代计算(见算法1中的第②~ ⑧步). Step2.寻找使得lI r埘ll最小的Y (见算法1中的第⑨ 步). Step3.计算Xm— 0+Vmy . Step4.如果残量不够小,增大Krylov子空间维度并重 复以上步骤. 本文采用的是重开始GMRES算法,即将Step4 改为:如果残量不够小,将 的初值置为 并重复 以上步骤. 算法1.GMRES算法 。 每次迭代的具体步骤 ①计算ro=6一A 。,卢一l Ir0 lI和 一,。/卢; ②初始化(m+1)×m阶的上Hessenberg矩阵H 为 0; ③循环』一1,2,…,m ④ 计算 一A ; ⑤ 循环 一1,2,…, ①http://en.wikipedia.org/wikilGeneralized-minimal_residual_method 第4期 柳有权,等:大规模稀疏线性方程组的GMRES ̄GPU快速求解算法 ⑥ 计算h 一 ・v ;计算 —wJ—hijv ⑦ 计算h 1.J=l1 ll,如果h 一0,则将 赋给 m,转到⑨; ⑧ 计算 + — /h ⑨计算l1日 Y 一 ll的最小值,得到对应的 ; ⑩计算 一X。+V . 设系数矩阵A为n阶可逆矩阵,每行非零元数 为k,Krylov子空间维度为m(常数),则每次迭代中 各种计算所需次数及复杂度如表1所示. 表1 算法1中各计算步复杂度分析 很显然,无论从计算量还是从GPU程序的优 化难度来说,表1中的前2项都是整个算法中最耗 时的部分. 设稀疏矩阵与矢量乘积、内积与范数的浮点运 算次数分别为厂 和厂 ,则f 一2nk( +1),f 一 er(m+3)+2 .当忌》 一予+1时,稀 疏矩阵与矢量乘积运算决定程序的整体性能;当 忌《 +1时,矢量内积与范数运算决定程序的整体 性能. 考虑到矢量内积和范数运算在GPU中运行比 稀疏矩阵和矢量乘积遇到的性能瓶颈要少很多,因 此可以认为m值越大,加速效果越好;k值越大,加 速效果越差. 尽管m较大时可以加速收敛和提高精度,但同 时会使每次迭代的计算量增加并且使存储占用增 加;另外由于浮点精度有限,因此m的值不宜取得 过大.至于m的值取多少合适,要考虑矩阵本身的 特点、机器内存的、浮点精度的等,一般根 据矩阵条件数凭经验决定.另外,过小的m可能导 致收敛停滞,文献E4]对此进行了一些讨论. 2 GMRES算法的CUDA优化实现 关于CUDA编程的详细说明可以参考NVIDIA 手册_7 ,本文不再赘述.根据本文前面的描述,GMRES 算法跟其他线性方程组迭代法求解类似.整个计算 被分解成稀疏矩阵与矢量乘积运算,矢量内积与范 数,稠密矩阵与矢量乘积,矢量加减、矢量乘标量,最 小二乘求解等几步.为了获得整体优化的效果,本文 对于不同的计算步骤采用不同的策略. 2.1 稀疏矩阵与矢量乘积运算 本文中矩阵存储使用压缩稀疏行存储(compressed sparse row,CSR)格式,即将所有非零元素存入一 个数组data,这些非零元素的列号存人另外一个数 组indices,而数组rpos存放每一行第一个非零元素 在数组data中的序号,rpos中的最后一个元素存 放总的非零元素的个数.CSR格式稀疏矩阵与矢量 乘积在GPU中计算遇到的瓶颈以及一些算法实现 见文献[11].本文改进了翟艳堂等的稀疏矩阵与矢 量乘积运算方法①,对矩阵数据读取全部实现了合 并访问,取得了很好的加速效果. 乘法与加法分为两步的思想较为简单,其步骤 如下: Step1.Kernel1.计算每个非零元与相应标量的乘积,保 存结果至全局存储器,用rares表示. Step2.Kernel2.读取Kernell的结果,并将每行对应的 数据累加得到结果矢量. Kernel1较为简单,这里不再赘述,详细请参考 相关链接.对于Kernel2,使用纹理存储并不能解决 非合并数据读取大量访问延时造成的性能瓶颈. 本文提出合并访问的Kernel2算法(以block大 小为128为例). 算法2.稀疏矩阵与矢量乘积运算第2步核心 Step1.tidb' ̄-thread index in Block Step2.tidg ̄thread index in Grid Step3.sum" ̄--0.0 Step4.define shared array rposs[129]and mres~S ElzC Steps.rpos~s[tid一6]・ —r户0s[tid—g] Step6.if tidb一0 then rpos5[129]一rpos[tid—g+ 128] Step7.baseAddress ̄--rpos5 Eo] Step8.rares~5[tid一 一mr∞[baseAddress+tid~6j Step9.P一{z J rflos_s[tid_ ≤x<rpos一5[tid_6+1])n {38l baseAddress ̄x ̄baseAddress+128) Step10.if P一 goto Stepl3 ①http:// da _dn. t/c0nt 。 pr0/ idia_ h。 . p 。intid一52 556 计算机辅助设计与图形学学报 第23卷 Stepl 1.for all elements i in P 约为51个.假设第一个线程在block中的编号为tid, Stepl2. s m+一lnl'ess[ ] 则tid 32为其在warp块中的编号.当tid 32≤12 Stepl3.baseAddress ̄--bassAddress+128 时,这51个线程占用2个相邻warp块;当tid 32> Stepl4.if baseAddress<rpos一5 Dz93 goto Step8 12时,这51个线程占用3个相邻warp块.因此, Stepl5.result[tid—g]一5Mm. 其中,rpos—s为每个block分配的共享内存 Stepl2中SM的利用率期望值为 13_(shared memory),用于存放该block所处理的某些 32 51十,一19 51・一64 1 32 96一I一●….64. 行第一个非零元素在数组data中的序号;mres—S 类似地可以计算当k一5时,SM的利用率约为 也是每个block分配的共享内存,用于存放该block 0.77. 所处理的非零元与相应标量的乘积.通过共享内存, 另外,扩大mres—S的容量可以减少从Step8~ 可以大大加速数据访问,减少延迟. Stepl4的迭代次数,因此该算法效果很好.例如当 本文通过分析发现上述算法存在如下不足:假 mresS容量为512,k<4时,从Step8 ̄Stepl4的迭 设稀疏矩阵平均每行非零元素为k,则每个block中 代在大多数情况下只需要进行一次.也就是说,所有 1 0o 同时执行Stepl2的线程约为 ,对于是一10,实际 需要的数据被一次性读入共享内存,然后所有线程 工作的只有13个线程,远不足一个warp(一个 同时对共享内存中的数据进行操作. warp由32个线程组成).也就是说从Step8~ 对Step9求交集,可以根据{ Ia≤x<b)n{zI Stepl4的每一次迭代中,当程序执行到Stepl2时, f≤x<d}≠ 甘一(6≤c V口≥ )很容易地完成. 1 0 Streaming Muhiprocessors(SM)的利用率只有 . 本文与文献[11]中的稀疏矩阵和矢量乘积算法 进行了代码级性能比较,结果如表2所示,其中本文 更糟糕的是,Stepl2是上述算法计算量最多的一步. 所用的矩阵数据全部来自http://www.cise.uf1. 解决的办法是扩大mres—s的容量.例如mres—S edu/research/sparse/matrices. 的长度为512,对于是一10,同时执行Step12的线程数 表2稀疏矩阵和矢量乘积算法效率比较 由表2可以看出,本文实现的稀疏矩阵和矢量乘 况下将每个标量乘积累加起来.本文采用reduction ̄ 积算法要比NVIDIA提供的算法效率高,而且对于一 的思想,即将2个输入矢量划分成若干对小矢量,每 些矩阵,NⅥDIA提供的算法存在失效的情况. 个block负责计算一对小矢量的内积,这些小矢量 2.2矢量内积和范数运算 的内积被写到mapped memory中,由CPU负责将 算法中矢量范数采用欧几里德范数,由于它与矢 这些小矢量内积加起来得到最终结果. 量内积计算类似,因此本文只介绍矢量内积的计算. 内积计算的困难在于如何在保持一定并行度的情 第4期 柳有权,等:大规模稀疏线性方程组的GMRES-GPU快速求解算法 为了对不同问题规模都能得到最佳的warp块 装载量,本文并未固定小矢量的长度,而是在 block个数不大于512的情况下动态地决定小矢量 设输入矢量长度为n,block大小为b,block 个数为g,小矢量长度为m,则有如图1所示的判 定树. 的长度,因此得到的小矢量不多于512对. 图1矢量内积计算判定树 从图1中可以看出,利用该判定树得到的线程 和数据划分在任何时候都不会对warp块装载量产 生影响.当,2≥512×512时,尽管此时block的数量 误差舍入.例如,有一个长度为1 000 000的数组,每 个元素均为0.001,使用IEEE标准单精度顺序相加 得到的结果是991.141 541,和正确结果1 000相差 被为512,但由于每个SM只能同时装载2个 大小为512的block,所以这个是合理的,不会 对性能产生影响. 接近百分之一.这是由于浮点精度导致的舍入误差 造成的.使用reduction方法的情况要好很多,如算 法3代码所示,数组元素是两两相加的.当数组元素 的数量级差别不大时,使用reduction方法得到的结 果几乎是精确的;但当数组元素数量级相差较大时, 其结果是偏小的,如1.OEIO+1.OE一10—1.0El0.解 以block大小为128为例(图1判定树根的右 节点),kernel函数代码如下: 算法3.CUDA内积运算核心 globalvoid innerPro128(float*v,float*'‘,,float* 决加法误差的方法有Kahan summation算法①,但 其会造成性能上的大幅度下降,因此本文并未采用. 2.3稠密矩阵矢量乘积、矢量加减、矢量乘标量和 tl ̄tre¥){ int tid—threadldx.z; int tidingrid—blockIdx.z*blockDim.z+ 最小二乘求解 threadldx.z: sharedfloat r s[128]; 稠密矩阵矢量乘积、矢量加减和矢量乘标量问 题较为简单,本文采用CUDA实现.CUDA SDK中 提供了类似的例子代码,这里不在赘述.求解最小二 r s[tid]一v[tidingrid ̄*w[tid—in—grid]; syncthreads(); if(tid<64)rS[tid]+一r—S[tid+64]; 乘问题Y—arg rain I lHmy—t3e l 较为烦琐口 ],l 但从表1中可以看出,最小二乘问题的浮点运算次 数大约为4m。,计算量并不大,没有构成问题的瓶 syncthreads(); if(tid<32){ r s[tid]+一r_5[£ +32]; if(tid<16)rs[tid]+一r—s[£ +163; 颈,因此该计算在CPU端完成. if(tid<8)rs[tid]+一r__s[£ +8]; if(tid<4)rs[tid]+一r-s[f +4]; 3实验与结果讨论 本文算法在Visual Studio 2005+Window if(tid<2)rsEtid]+一r—s[£ +2]; if(tid ̄1)mres[blockIdx.z]一r一 [o]+r—sill; XP/Vista软件环境下,采用C++语言结合CUDA ) 3.1编程实现.GPU为NVIDIA GeForce GTX260, 分别在2台装有不同CPU的机器上运行,机器1为 其中 和W分别为输入的2个矢量,而mres则 为生成的小矢量. 矢量内积计算遇到的另一个问题是数组求和的 558 计算机辅助设计与图形学学报 第23卷 Intel Core 2 Quad CPU Q9400@2.66 GHz,得到了 CPU 920@2.67 GHz,得到平均20余倍的加速效 平均40余倍的加速效果,机器2为Intel Core i7 果.它们均迭代1O次,具体实验结果如表3所示. 表3机器1。2上的性能测试 0.234 0.563 21.407 0.672 61.782 3.062 32 208.O31 4.782 43.5 由于机器2的CPU配置比机器1高1倍,导致 从表4可以看出,曲线收敛速度跟矩阵关系比 机器2上的GMRES算法的CPU版本运行的效率 较紧密,有些收敛得快,有些收敛得慢,说明收敛效 比机器1的提升了2倍,但同样的GPU版本代码效 果首先取决于方程组系数矩阵的实际情况.对于同 率却略有下降,这可能是数据总线传输速率差异导 一个矩阵来说,Krylov子空间维度的大小则对于收 致的. 敛有着至关重要的影响.另外可以看出,rajat26和 为了确保GPU与CPU运算结果一致,本文进 c一26 2个矩阵的求解并没有收敛到近似解,这是矩 行了收敛曲线和误差分布分析,总共迭代100次,取 阵的条件数较大对误差较为敏感造成的,从收敛曲线 了5个矩阵样本;其中测试矩阵得到的收敛曲线和 也可看到这2个矩阵的计算收敛情况出现异常. 误差分步直方图如表4所示. 第4期 柳有权,等:大规模稀疏线性方程组的GMRES-GPU快速求解算法 559 表4收敛曲线与误差分步直方图 收敛曲线 (横坐标为迭代次数,0 纵坐标为lg__cP 子空间维度 ) 误差分步直方图 O (横坐标为标量误差I I。. 吨如 l,纵坐标为所占比例) 1.O O.8 l l l 0.6 0.4 l l 0.2 } \ 1 11 21 31 41 51 61 71 81 91 L I O ≥l_0 |_ J l l—L_lI_-IJ—■ 1.0E一2 1.013-4 1.0E一6≤1.0E一7 I- ■ I .1.O O.8 1 3 5 7 k—~ 、~ ——~—0.6 16 — —— 0.4 . l 9 | 、 \ _ O.2 l l ≥1.0 I _ 1 . 1 1.0E一2 1.0E一4 1.0E一6《1.0E一7 11 O 11 21 31 4l 51 61 71 81 91 L J l l - L一■.1 1 l l O.8 O.6 ’●L 、 1.O ~l 0.4 32 — ^^ N~. l 0.2 lO l I l L .—,-_、,_ ^ 一 l一 ≥1.0 . I l L _ 一 1.0E一2 1.013—4 1.0E一6《1.0E一7 1 11 21 31 41 51 61 71 81 91 注:——c-26;——bcsstkl5;-C-26: _bcsstkl5: ——crystk02;——rajat26;——ASIC_680ks; ’crystk02; -rajat26; -ASIC_680ks 我们拟将本文中的求解算法应用到流体仿真和变形 4结论和未来工作 本文通过对GMRES算法进行算法分析,并利 用CUDA实现线性方程组的快速求解.尤其是对于 稀疏矩阵与矢量乘积运算部分,通过合并访问和共 享内存策略提高数据负载的利用率,使得对于大规 模线性方程组的求解效率得到提升,对于当前主流 的PC机来说,GPU版本要比CPU版本快20~40 倍,对于某些矩阵效率提升更多.由于GMRES迭 计算的计算机动画研究中去. 致谢感谢NVIDIA公司提供GTX260显卡, 感谢佛罗里达大学提供大规模稀疏矩阵测试数据! 参考文献(References): Saad Y,[1] Schuhz M H.GMRES:a generalized minimal residual algorithm for solving nonsymmetric linear systems [J].SIAM Journal on Scientific and Statistical Computing, 1986,7(3):856-869 代算法适用于一般的线性方程组的求解,所以只要 求系数矩阵可逆就有解.虽然条件数对迭代收敛结 erative methods for sparse linear systems[M].2nd E2] Saad Y.Ited.Philadelphia:SIAM,2003 果会有影响,相信在工程应用领域会大有用武之地. 56O 计算机辅助设计与图形学学报 第23卷 1-33 Quan Zhong,Xiang Shuhuang.A GMRES based polynomial prec0nditi0ning algorithm[J].Mathematica Numerica Sinica, 2006,28(4):365—376(in Chinese) (全忠,向淑晃.基于GMRES的多项式预处理广义极小残 差法[J3.计算数学,2006,28(4):365—376) [4]Habu M,Nodera T.GMRES(m)algorithm with changing the restart cycle adaptively Ec]//Proceedings of Algoritmy Conference on Scientific Computing.Heidelberg:Springer, 2000:254-263 [53 Wu Enhua,Liu Youquan.General purpose computation on GPU[J].Journal of Computer—Aided Design&Computer Graphics,2004,16(5):601—612(in Chinese) (吴恩华,柳有权.基于图形处理器(GPU)的通用计算[J]. 计算机辅助设计与图形学学报,2004,16(5):601—612) [6]Wu E H,Liu Y Q.Emerging technology about GPGPU [C]//Proceedings of IEEE Asia Pacific Conference on Circuits and Systems.Los Alamitos:IEEE Computer Society Press, 2008:618—622 [7]NVIDIA CUDA C programming guide.Version 3.1[M]. San Jose:NVIDIA,2010 [8]Wang M L,Klie H,Parashar M,et a1.Solving sparse linear systems on NVIDIA tesla GPUs[M]//Lecture Notes in Computer Science.Heidelberg:Springer,2009,5544:864— 873 [9]Velamparambil S,MacKinnon-Cormier S,Perry J,et a1. GPU accelerated Krylov subspace methods for computational electromagnetics[c]//Proceedings of the 38th European Microwave Conference.Los Alamitos:IEEE Computer Society Press,2008:1312—1314 [10]Ghaemian N,Abdollahzadeh A,Heinemann Z,et a1. Accelerating the GMRES iterative linear solver of an oil reservoir simulator using the multi—processing power of compute unified device architecture of graphics cards[c]// Proceedings of the 9th International Workshop on State—of—the—Art in Scientific and Parallel Computing. Heidelberg:Springer,2008:156—159 [11]Bell N, Garland M. Efficient sparse matrix—vector multiplication on CUDA JR]. San Jose: NVIDIA, NVR一2008—004。2008 [12]NVIDIA CUDA C best practices guide.Version 3.1[M] San Jose:NVIDIA,2010