CUDA简介

CUDA(Compute Unified Device Architecture),是显卡厂商NVIDIA推出的运算平台。 CUDA™是一种由NVIDIA推出的通用并行计算架构,该架构使GPU能够解决复杂的计算问题。 它包含了CUDA指令集架构(ISA)以及GPU内部的并行计算引擎。 开发人员现在可以使用C语言来为CUDA™架构编写程序,C语言是应用最广泛的一种高级编程语言。所编写出的程序于是就可以在支持CUDA™的处理器上以超高性能运行。CUDA3.0已经开始支持C++和FORTRAN。

GPU硬件结构

GPU硬件架构和CPU硬件架构有着本质上的不同,图2-5显示了一个位于PCI-E总线另一侧的多GPU系统。 在下图中我们可以看到PCI-E线连接多个GPU,然后图中展开一个GPU,这里值得注意的是,一个GPU其实就是一个SM(Stream Multiprocessor)阵列,这里具体有多少个SM有显卡决定,不同显卡数量不同。每个SM中又有多个SP(Stream Processor),这里的SM就是平时所说众核处理器中的‘核’了,而SP就是嵌在每个核中的计算核心。 block(线程块,2.2.3线程模型中将具体介绍)是执行在SM上的, SP数量由具体显卡决定,每个SM中SP数量由显卡架构决定,所以如果显卡架构相同,SP数量越多SM数也就越多,显卡真正的并行能力是由SM数量和每个SM中的最大线程数决定,例如有一块显卡有2个SM,而每个SM的最大线程数是10,那么物理上真正同时执行的线程数量就只有20,如过你开了40个线程,那么理论上执行时间是你开20个线程的两倍。

GPU architechture

CUDA线程模型

由于CPU和GPU设计上的不同,CPU每个核心有着很强的计算能力,但是核心数量有限,通常只有两个或者四个,所以CPU适合处理少量复杂的任务,而GPU则相反,GPU有着大量的SP,但每个SP的运算能力不如CPU,所以GPU适合解决大量简单的任务。从存储器的角度来看,CPU只有少量的寄存器,用来线程上下文切换昂贵,因为切换进程时,必须要将寄存器中的内容保存至内存,恢复时又要从内存写会到寄存器,这也正符合CPU设计的目的,适合处理少量复杂任务,不适合线程大规模的并行。GPU拥有众多寄存器组,线程切换非常简单只要进行寄存器组调度就可以了,所以非常适合多线程并行程序。

线程是并行程序的基础构建模块,线程束(warp)是GPU线程执行的基本单位,在我的显卡上大小为32,就是说一个线程束由32个线程组成,不同显卡可能会有不同。GPU是一组SIMD向量处理器的集合,每一个线程束中的线程一起执行,理想情况下一个指令只需要读取一次,然后广播给同一个线程束中的线程。CUDA的线程模型分为两层粗粒度的块并行和细粒度的块内线程并行。

  • 细粒度并行:细粒度并行是指块内的线程并行,块内线程可以通过共享内存通信,CUDA中的线程同步命令__syncthreads()也只对同一个块内的线程有效,CUDA允许你指定块内线程的布局,例如每行36个线程,两行还是每行16个线程4行,GPU和CPU一样是行优先顺序读取连续内存的。

  • 粗粒度并行:粗粒度并行是指块间并行,块间可以通过全局内存通信,线程组成线程块,而线程块组成线程网格,CUDA允许你制定线程网格的布局例如每行2个块,两行或者每行1个块,4行。下图便是一个线程模型样例。

thread model

CUDA内存模型

GPU内存通常分为以下几种,寄存器,共享内存,纹理内存,常量内存和全局内存。寄存器是访问速度最快的,但数量有限,每个SM只有有限的寄存器,如果你每个线程使用寄存器过多,这将最终限制可同时执行的线程数,并且申请内存超过寄存器最大数量后,多余的数据将会被存储到LocalMemory,它的读写速度和全局内存一样慢。共享内存的访问速度仅此于寄存器,用户可以通过共享内存实现块内线程的通信,通常你可以将块内所有线程都用到的数据放到共享内存中。常量内存是全局内存的一种映射,是只读内存(对于GPU而言),相较于全局内存,常量内存有两个优势,其一是常量内存有高速缓存,也就是说当连续访问相同的常量内存时,可以只用从常量内存读取一次。第二点是常量内存的读取会进行warp内的广播,同一个warp只需读取一次。和常量内存一样,纹理内存也是全局内存的映射和只读(对于GPU)内存,在特定的访问模式中,纹理内存同样能够提升性能并减少内存流量。纹理内存缓存在芯片上,因此在某些情况中,它能够减少对内存的请求并提供更高效的内存带宽。纹理缓存是专门为那些在内存访问模式中存在大量空间局部性(Spatial Locality)的图形应用程序而设计的。全局内存的访问没有任何限制,内存大小也是最大的,相应的缺点是内存访问最慢且没有硬件优化(如高速缓存)。下表展示了各种内存的带宽和访问延迟。

内存访问速率

常见优化方法

  • 每个流处理器簇(CUDA详细介绍见2.2)都有线程上限限制分配的线程数,例如本人的显卡(Nvidia GT 525M)每个SM最多分配线程数为1536,这时假设我的显卡有两个SM,我开了3072个线程,我们现在将这些线程分配到3个线程块,每个线程块分配到1024个线程由于每个SM无法分配到两个线程块,所以就需要3个SM,这就导致一次并行只能执行2048个线程没有充分发挥显卡的并行性,如果我们分配两个线程块,每个线程块分配1536个线程,则所有线程可以一起并发执行。

  • 显卡的执行单位是warp,warp大小为32,如果线程数量有一定灵活性的话,最好分配32倍数的线程数量,而在程序算法实现中,应该尽量减少分支,因为GPU的分支执行逻辑与CPU不同,CPU会利用空间局部性原则和时间局部性原则进行预测,命中率决定效率,而GPU相比下则简单的多,一个warp的线程一起执行一个分支,假设有10个线程为true,进入分支,而剩下22个为false的不会继续执行下面,他会被设为未激活状态然后等待,直到那10个线程执行完,然后遇到另一个分支,这时候状态和刚才相反,10个线程变为未激活状态等待,另外22个线程执行分支,也就是说如果一个wrap中的线程在这种情况下执行时间是所有分支执行时间之和。

  • 在循环中尽量少的进行函数调用,在循环中尽可能的展开你的函数,这将能大幅提高你的kernel执行效率。

  • 在计算密集的kernel中尽量避免整数除法和求模的使用,如果可以人为求出就人为求出,不可以的话可以尝试通过移位等方法替代。尽量少的使用浮点数除法,通常可以通过调整计算顺序来达到此目的。


ShiweyYan

A game developer who graduates from SCUT.