并行计算课程算法实践教学的新工具CUDA编程模型.docx
《并行计算课程算法实践教学的新工具CUDA编程模型.docx》由会员分享,可在线阅读,更多相关《并行计算课程算法实践教学的新工具CUDA编程模型.docx(14页珍藏版)》请在冰豆网上搜索。
并行计算课程算法实践教学的新工具CUDA编程模型
》》》一
一一
puterEducation
文章编号:
1672-5913(2008)23—0103-04
“并行计算"课程算法实践教学的新工具:
CUDA编程模型
王智广1,刘伟峰2
(1.中国石油大学(北隶)计算机科学与技术系,北京102249;
2.中国石油化工股份有限公司石油勘探开发研究院信息技术研究所,北京10008
3)
摘要:
本文首先从当今多核微处理器的发展趋势出发,介绍适用于多核微处理器的细粒度并行编程模型CUDA,以及其适用于“并行计算”课程教学的一系列优势,接着对当前可用的几种并行编程模型进行分析和比较,最后给出采用CUDA编程模型的一个矩阵与矩阵相乘的教学实例。
关键词:
并行计算;实践教学;CUDA
中图分类号:
G642
文献标识码:
B
1引言
1998年,教育部高等学校计算机科学与技术教学指导委员会将“并行计算”课程定位在高等学校计算机专业高年级本科生或研究生以及面向计算学科的非计算机专业的研究生层次上。
“并行计算”课程一般主要包括以下四个层面:
(1)在算法理论层面主要讲授可计算性与计算复杂性以及算法研究的数学基础;(2)在体系结构层面主要讲授并行计算机体系结构;(3)在算法设计层面主要讲授计算机科学及其他科学计算领域中的非数值和数值并行算法;(4)在算法实践层面主要讲授基于并行编程模型的并行算法编程。
实际上,对于大多数学习“并行计算”课程的非并行计算研究方向的学生来讲,第(4)层面,即算法实践是课程最具价值的组成部分。
将并行算法实现成程序并投入大规模科学与工程计算生产应用是课程的最主要学习目的之一。
其中对于编程模型的选择又是算法实践教学的重要前提。
以往“并行计算”课程在算法实践教学中大多采用MPI作为编程模型,可MPI的粗粒度特性不仅不足以满足“并行计算”课程的实践需要,还很难利用当前多核微处理器的硬件细粒度优势。
本文从当今多核微处理器的发展趋势出发,介绍适用于多核微处理器的编程模型CUDA,以及其适用于“并行
计算”课程教学的一系列优势,并给出对几种主流并行编程模型与CUDA的分析和比较,最后提供一个基于CUDA的矩阵与矩阵相乘程序作为教学实例。
2多核微处理器发展趋势
过去的20年间,提高运行频率是CPU性能提高的主要方式,然而,自2003年以来,这种趋势发生了变化,不断提高的CPU频率带来了高功耗和高发热量,使得主流CPU频率止步于4GHz,并向单芯片多处理器(ChipMultiProcessors,CMP),即多核方向发展。
2005年,Intel和AMD正式向主流消费级市场推出了双核心的CPU产品,2007年推出了4核心的CPU,按照各厂商的发展路线图,今后大约每2年单CPU上的核心数将翻番。
但由于市场变化和研制成本的原因,多核CPU的每个核都基于以往CPU的单核设计,保留了如乱序执行等很多单核时代的复杂执行方式,使得其对于科学计算等问题的计算能力提高非常有限。
此时,以游戏加速和图形处理为初衷设计的GPU(graphics
processing
unit,图形处理器)以超出摩尔定
律的速度发展,并开始在非图形的高性能计算领域被大量使用。
图1为近年来同等市场价格CPU和GPU的浮点运算速度对比图。
作者简介:
王智广(1964-),男,教授,CCF高级会员,主要研究方向为分布式计算和并行计算;刘伟峰(1981一),男,中国石油大学(北京)计算机科学与技术系2006届硕士,工程师,CCF会员,主要研究方向为图形处理器架构,高性能计算可视化。
万方数据
”卜…一1而磊蕊i墨圈隧蕊————…
镶镧《
啪乏享J
啪““m
/
||
∞乱。
一Lo
抛e-。
。
=/。
71/
88罗/厂
30GH:
御
/
Core2Quad
3oGHz
w!
:
2^磊蒿:
意’。
.——/
o
JanJuIJanJulJanJuIJanJuIJanJuI2003
2003
2004
2004
2005
2005
2006
2006
2007
2007
图1
GPU和CPU的FLOPS理论峰值
由图可见,GPU的浮点运算速度可以达到CPU的5倍以上。
带来这种数据处理能力差别的主要原因在于GPU最早为并行处理大量三维计算机图形学中的顶点和像素数据而设计,近年来为通用计算又进行了一系列改进。
其天生并行的体系结构决定了GPU非常擅长以并行的方式运行高运算强度的应用。
以nVidiaGPU为例,与CPU体系结构相比,GPU体系结构的优势主要有:
(1)单芯片上的更多处理器(核)与G80核心GPU包含的128个核心相比,CPU目前最高仅达到4核,虽然CPU每个核的运算能力高于GPU上的单核,但后者凭借更多核的并行在计算能力上高出很多;
(2)将更多的晶体管用于计算单元GPU运算能力远高于CPU的重要原因是GPU将更多的晶体管用于计算单元。
目前CPU将约70%的晶体管用于Cache;而GPU晶体管的80%以上是计算单元;
(3)超多活动线程G80GPU能够保持12288个活动线程的上下文,与此相比,CPU通过线程池管理的线程不过数十个。
GPU超多活动线程可以隐藏大量计算单元的访存延迟,在处理数据敏感的应用时要比CPU的Cache方式更加有效;
(4)高显存带宽
GPU与其显示设备内存间的带宽超
过CPU与内存的带宽10倍以上。
配合超多活动线程,使GPU十分适合大规模并行处理高运算强度的应用。
正是由于这些原因,在国际高性能计算研究领域,针对以GPU为代表的硬件细粒度并行计算的研究是近年来的一个新的热点。
并诞生了大量基于GPU的科学计算改进算法,其领域涵盖各类数值模拟方法,包括计算天体物理、计算流体力学、计算量子化学,甚至计算金融学在内的众多领域已经从GPU计算中获益。
万方数据
由此可见,相对于多核CPU而言,GPU的高度并行
体系结构天然地适合进行并行计算的研究和教学。
而在2007年以前,GPU的并行程序设计还主要是通过OpenGL或Direct3D这种图形API来完成,入门难度较高。
所以在高性能计算的研究者们对于利用GPU进行并行计算的迫切需求下,CUDA编程模型诞生了。
2
CUDA编程模型
CUDA(Compute
UnifiedDevice
Architecture,统一计算
设备架构)是2007年由nVidia推出的一套并行编程模型。
其推出的最初目的是为解决GPU通用计算的易用性问题。
至今已有众多研究者利用GPU的高度并行性特点将科学计算算法迁移至CUDA编程模型并在GPU上获得了相对于CPU平均数十倍的性能提升。
随着CUDA编程模型被越来越多的研究者接受,
UniversityofIllinoisat
Urbana-Champaign的研究人员利用
源.源的解决方案设计了MCUDA(Multicore.CUDA),使CUDA程序也可以运行在多核CPU上,且运行效率要高于多数原生的CPU优化方法。
使CUDA的应用范围基本覆盖了包括CPU和GPU在内的当今主流多核微处理器。
可以说,CUDA并行编程模型的出现和迅速发展标志着细粒度并行程序设计已经逐渐为主流程序设计人员所接受。
CUDA编程模型有SIMT和显式数据调用2个主要特点。
2.1
SIMT运行方式
与以往SIMD和MIMD的并行计算常用指令和数据流
运行方式相比,CUDA程序的运行方式被称为SIMT(Single
InstructionMultiple
Thread,单指令多线程),其意义是可以
让程序设计人员把编程的SIMD硬件对象作为若干个标量处理器调用,即SIMT没有固定的矢量宽度,并允许每一条线程占据各自路径,在不受SIMD可编程性制约的前提下保证每个处理器核心一直处于完全被利用状态。
这样在并行程序设计上就可以将精力专注于程序线程的拆分,达到提高并行化效果并且降低程序编写难度的目的。
2.2显式数据调用
CUDA的访存指令与C语言相同,但增加了一项重要特性,即对共享内存(SharedMemory)的显式操作。
共享内存是GPU芯片内的一块分为16个bank的16KB存储器,主要作用是作为并发线程间的共享数据存放空间,另外也可以显式地将其当成Cache使用,存放并发线程其需要多
》》卜—————]囊豳豳霹里里型塑型!
!
一
斓镧镧
次存取的数据。
当每个线程访问不同bank的时候,共享内存的存取延迟与寄存器的存取延迟一样低,即数百倍地低于片外显示内存或主存的存取延迟,所以对共享内存的有效显式操作是提高CUDA并行程序运行效率的关键。
本文第4节会给出一个使用共享内存的教学实例。
总体来讲,CUDA编程模型中基于多核的SIMT细粒度并行方式和对存储器的显式操作是十分理想的课程教学工具,有助于学生设计既贴近硬件体系结构又独立于运行平台的并行程序。
虽然学生在算法实践初期会觉得细粒度的并行程序设计较难接受,但就我们的教学经验来看,具有计算机体系结构和编程语言基础的学生大约在4周内可以对CUDA的并行编程模型基本熟悉,况且在算法实践中获得数十倍性能提升带来的成就感也会促使学生逐渐克服学习的困难。
3
几种弗行编程模型的分析和比较
目前国内外一般教材和课程教学选用的并行编程模
型主要有3种:
消息传递接口MPI、线程接口POSIX和基于指令的OpenMP。
这三种编程模型各有比较鲜明的特点。
本节将它们和本文选用的CUDA编程模型一并进行分析和比较。
3・1
POSIX
POSIX(PortableOperatingSystemInterfaceofUnix,可
移植Unix操作系统接口)线程接口也称Pthreads,1995年被IEEE通过成为被绝大多数厂商支持的线程API。
它提供了处理诸如死锁和竞态条件这类问题的工具,但没有限定线程的具体工作方式,对于如何编制线程并行的程序留有很大的余地。
基于此,Pthreads被认为过于底层和编程难度较高,所以目前多数“并行计算”课程不将Pthreads作为主要的算法实践工具。
3.2
OpenMP
作为共享存储标准而在1997年问世的OpenMP是为在多处理机上编写可移植的多线程应用程序而设计的一个应用编程接口。
OpenMP编程模型包括一套平台无关的编译指导(pragmas)、编译命令(directive)和一个用来支持它们的函数库。
OpenMP显示地指导编译器如何利用应用程序的并行性,而开发人员不需要关心实现细节,这使得OpenMP的学习入门难度相对较低,但同时也很难完成并
万方数据
行算法的课程教学要求,所以一般也不将OpenMP作为主要的算法实践工具。
3.3
MPI
1994年由政府、学术界、产业界共同制定的MPI(MessagePassingInterface,消息传递接口)是根据并行应用程序对于消息传递的需求而定义的一组标准接口说明和不同厂商提供的相应方法实现。
MPI吸收了PVM等众多消息传递模型的优点,在集群型高性能计算机的流行化趋势下成为了目前最流行的并行编程模型。
虽说MPI的学习曲线较长,但大多数“并行计算”教科书还是主要使用MPI作为并行算法教学实践的工具。
3.4
4种并行编程模型比较
表1是4种并行编程模型主要特性比较表。
表1
4种并行编程模型主要特性比较表
POSIX
OpenMPⅣ衅ICUDA
特征
低级原语共享存储消息传递单指令多线程并行粒度细粒度细粒度粗粒度细粒度存储模式共享存储共享存储分布存储共享存储数据调用方式显式隐式显式显式学习难度难容易较难一般可扩展性
较好
较差
好
较好
4矩阵与矩阵相乘并行算法教学实例
矩阵与矩阵相乘并行算法是最常用也最具代表性的并行算法之一,十分适合作为学生对并行算法实践的入门
教学。
本节描述一个基于CUDA编程模型和块矩阵算法的矩阵与矩阵相乘并行算法教学实例。
4.1串行算法
对于两个胛×"稠密矩阵4与B相乘得到乘积矩阵C=A×B的算法,易写出其串行实现的伪代码如下:
procedureSERIAL—MAT—MULT(A,B,C)
Begin
fori:
=0ton一1do
for
J:
=0
to
n一1
do
Begin
C[i,j]:
=0j
for
k:
=0
to
n一1
do
C[i,J]
:
=C[i,J]
+A[i,k]
×
B[k,J];
endforj
endSER工AL
MAT
MULT
¨努—弋而蕊蕊i暖圈酲峨
4.2块算法
针对粗粒度的并行机制,一般采用基于块矩阵运算矩
BeginCsub:
=0;for
k:
=0
to
q—IdoASi?
k;BSk?
j;
j
斓碉镧
阵与矩阵相乘算法,即将整个矩阵分成矩阵块大小为(哟)
X(n/O的块矩阵,把原矩阵的代数运算转换成对这些块矩阵中元素的代数运算,进而通过消息传递实现粗粒度的并行。
块算法的伪代码如下:
procedureBeginfor
i:
=0
to
q一1to
do
do
loadload
Ai?
ktoBk?
J
to
syncthreadsfor
i:
=O
to
15do
csub:
=Csub
+A鼠,k[threadIdx.x,1]×BSk,j[_z,threadIdx.y];
BLOCK—MAT—MULT(A,
B,
C)
syncthreads;
endfor;
Q,j[threadIdx.x,threadIdx.y]=Cs。
b;
endCUDA
MAT
MULT
for
J:
=0q一1
Begin
q,j:
2
for
0j
to
q一1
do
其@
syncthreads;语句表示在开始下一步操作之前
k:
=0
对线程块内的线程进行同步。
threadldx.x和threadldx.Y是二维线程块中线程的编号,通过它们也容易理解CUDA编程模型的单指令多线程的工作方式。
同时可以算出,为计算每个块G,,若不使用共享内存,
ei,j._Ci,J+At,k。
Bk,J;
endfor;
endBLOCKMAT
MULT
4.3块算法的CUDA并行实现
CUDA的细粒度并行机制要求从每个线程的角度考虑整个矩阵相乘,同时CUDA的grid-block-thread三级线程管理结构也要求对线程进行适当的块划分,本实例中对线程块的划分可以与块矩阵运算思想直接对应起来。
首先,根据需要划分两个胛×聍稠密矩阵彳与B为口
X
需要访问存放块矩阵Aj,t和取,的全局内存2n(响)2次,反
之仅需访存2q(,吻)2次。
通过以这种方式分块计算,可以有效利用快速的共享内存,进而节省许多全局内存带宽。
由于篇幅所限,本实例仅包括块矩阵相乘的CUDA代码主体部分,对代码的调用方式等其他内容请参考CUDA编程指南。
q块大小为(毗r)×(确r)的块矩阵,根据GPU的体系结构,
5结束语
本文在介绍多核微处理器发展趋势和CUDA编程模型的基础上,对CUDA和以往其他3种并行编程模型进行了比较,最后给出了一个矩阵与矩阵相乘的教学实例对CUDA编程模型进行进一步详述。
能够看出,CUDA作为近两年新出现的并行编程模型可以很好地利用新一代CPU及GPU芯片上数量众多的核进行细粒度并行计算,再结合CUDA相对易于学习掌握的特点,可以得出结论,CUDA编程模型可以与粗粒度的MPI互为补充,为学生带来更丰富的算法实践经验,并且完全可以作为“并行计算”课程算法实践教学的工具使用。
圈
取n/q=16以便每线程块的线程数是warp(GPU上32个线程组成的线程簇)大小的倍数,且低于每线程块的最大线程数。
然后,计算c的每个块Cf,,其执行过程是:
(1)使用每线程加载块矩阵4抽和圾,的一个元素,将爿城和瞰,从全局内存加载到共享内存;
(2)每个线程计算结果块矩阵的一个元素,其中每个乘积的结果累计到寄存器中:
(3)执行完毕后将寄存器中的结果写入全局内存中块矩阵Cf,的相应位置。
求每个块矩阵C“的CUDA伪代码如下:
procedureCUDA
MAT
MULT(A,B,C)
参考文献
[1】陈国良,孙广中,徐云,吴俊敏.并行计算课程的教学方法[J】.中国大学教学,2004,(2).[2】陈国良.并行计算一结构・算法・编程(修订版)[M].北京:
高等教育出版社,2003.[3】刘伟峰,王智广.细粒度并行计算编程模型研究[J].微电子学与计算机,2008,(10).[4】陈国良,安虹,陈岐,郑启龙,单久龙.并行算法实践[M].北京:
高等教育出版社,2003.[5]张林波,迟学斌,莫则尧,李若.并行计算导论【M].北京:
清华大学出版社,2006.
[6]nVidia.CUDAProgrammingGuide【EB/OL].http:
//www.nvidia.com/cuda.
万方数据