CUDA初步介绍

写在前面

最近由于项目的需要,我开始学习了GPU编程的一个框架–CUDA。CUDA由Nvidia公司开发,在并行计算领域获得了较大规模的使用。我这里总结下最近学习的一部分理论知识。

GPU发展历程,为什么我们选择GPU编程

通常意义上我们所使用的计算机架构是冯诺依曼架构,但是近些年来冯诺依曼架构的计算机遇到了明显的瓶颈。

存储墙

冯诺依曼系统架构简单来京就是处理器从存储器中不断取址、译码、执行指令,但是如今这种系统架构遇到了瓶颈:内存的读写速度跟不上CPU时钟频率。具有此特性的系统被称为内存受限系统,为了解决这样的问题,传统的解决方法是使用cache技术。通过给CPU设立多级缓存,可以大大的降低存储系统的能力。作者之前看到的处理器放大图上可以看出,CPU芯片上很大的一部分面积都是被各级缓存占据了。

但是很多事都是物极必反,随着缓存容量的增加,使用更大缓存所带来的收益会迅速下降,这也意味着需要寻找新的办法了。这体现了存储器延迟很难降低,缓存占据了芯片上的大部分面积(>70%),这也就是常常可以听到的Memory Wall。

功耗墙

由于芯片的集成水平越来越高,在很小的额空间里集成的电路和器件越来越多,产生的热量也越来越大(实际上CPU器件的产热除了和器件的集成程度相关之外大致上和CPU频率的3次方成正比),这也就以为着在CPU工作时要么不能有效的提高工作频率,要么不能合理的控制CPU温度。

物理极限

以现在最快的发展速度来看,芯片电路可以叨叨2-3个纳米级别,然而在这个级别上只能容纳10个原子–受限于量子的不确定性,晶体管将变得不可靠,也就是说摩尔定律在我们现今的硅系器件上失效了。现在的解决方案是寻找可以替代硅片技术的新材料。

多核解决方向

针对上面所提到的这些问题,芯片制造商对内部电路重新进行了设计,采用了多核处理器的方案,但是这也就意味着一个问题要分解成多个部分,并且采用并行算法设计程序,但是有的时候这样做的难度是很大的。

多核和众核时代

多个适当复杂度,低功耗核并行工作
时钟频率基本不变

  • 未来单核处理器不会更快,但计算机硬件会更宽
  • 必须重新设计算法

GPU简介

GPU的全程是Graphic Processing Unit,即专用图像显示设备。主要有以下几种:

独立显卡

主要的制造商是NVIDIA和ATI(Intel的Larrabee),是一个独立的硬件设备,主要的形态是插在PcI-E线上的显卡设备。

核心显卡

集成在CPU内部的显卡,与CPU融合更好。
AMD APU (加速处理器)
INTEL sandy bridge(SNB)、ivy bridge(IVB)

主板集成显卡

现在已经很少见了,这种显卡使用的是系统的主内存作为显存。近些年来我只在小霸王的游戏机上见过这种设计方式。

GPU的产生

GPU 产生之前,处理 2D 、 3D 图像都是依赖于 CPU ,但是由于 CPU 任务繁多,而且还有设计上的原因,这样面对日益复杂的 3D 图形图像时就会常常出现显卡等待 CPU 数据的情况,正是在这种情况下一种全新的图形图像处理器诞生了,从而大大加快了图形图像的处理速度,这种处理器就是 GPU。

传统的GPU

没有通用计算能力,主要用于到显示设备的图像显示:

GPGPU

后来由于人们发现GPU上可以实现数据级并行,解决计算一致性,并且GPU使用专用的存储器通道可以有效的隐藏存储器延时,也就发展了基于GPU的通用计算方法:General Purpose Computing on GPU

这种计算方式的核心思想就是用图形语言描述通用计算问题,并把数据映射到vertex或者fragment处理器。但是这样的处理方式也有显著的缺点就是硬件资源利用不充分,存储器访问方式严重受限,难以调试和查错,并需要掌握高级的图像处理知识和编程技巧。
由GPGPU演化出的重要计算模式,就是在异构协同处理器计算模型中将CPU和GPU结合起来加以利用。应用程序的穿行部分在CPU上运行,而计算繁重的部分则由GPU的高性能计算部分进行,而从用户的角度看,应用程序只是运行的更快了,获得了更好的性能提升。

这里提到的异构协同处理器的计算模式可以在很多高性能计算机上找到,并且如果我们换个思路的话,我们也可以由小见大,GPU又何尝不是一个微缩的超级计算机呢,毕竟GPU可以从上世纪的CELL机等当时的超级计算机身上找到影子,尤其是CELL机上的POWER-PC和流式多处理器的架构和任务分发机制与今日的GPU架构有很多的相同点。

并行算法设计流程

  • 任务划分 Partitioning
  • 通信 Communication
  • 任务组合 Agglomeration
  • 处理器映射 Mapping

要点:先尽量开发算法的并发性和扩展性,其次考虑通信成本和局部性,再次利用局部性相互组合减少通信成本,最后将组合后的任务分配到各个处理器。

## 现有的并行编程模型
分布存储编程模型:MPI
数据并行(数组划分)编程模型:HDF
共享存储编程模型:OpenMP、pthreads、TBB
专用异构编程模型:专用于NVIDIA GPU开发的 CUDA(Compute Unified Device Architecture)、开发MIC程序使用的LEO、Microsoft公司的Direct Compute等
通用异构编程模型:OpenCL(Open Computing Language)、OpenACC、OPENMP4.0
图形学API编程:利用DirectX和OpenGL等图形API进行程序映射,将需要计算的科学问题转换为图形处理问题,然后调用相应的图形处理接口完成计算。其后出现了相对高级的着色器语言,比如基于DirectX的HLSL和基于OpenGL的GLSL,以及同时支持两者的Cg

对现在比较常见的异构处理器编程模型来讲,代码量OpenACC < CUDA C <OpenCL,并且可以在GPU集群中方便的使用CUDA+MPI+OpenMP进行混合编程。

GPU与CPU的对比

  • CPU: 强控制弱计算,更多的资源用于缓存
  • GPU:强计算弱控制,更多资源用于数据计算

    在CPU+GPU的异构系统中,CPU与GPU经北桥芯片通过AGP或PCI
    -E高速总线相连,各自有独立的存储器。
    CPU与GPU间的通信是必不可少的,由于PCI-E的通信带宽限制,优化数据通信开销是必须要考虑的问题,也是影响程序性能的关键因素。

CUDA软件

CUDA(Compute Unified Device Architecture,统一计算设备架构)是由 NVIDIA 推出的通用并行计算架构,该架构使 GPU 能够解决复杂的计算问
题。 它包含了 CUDA 指令集架构(ISA)以及 GPU 内部的并行计算引擎。
CUDA 是一个全新的软硬件架构,可将 GPU 视为一个并行数据计算的设备,对所进行的计算进行分配和管理,无需将其映射到图形 API。

CUDA软件环境

CUDA的基本思想是支持大量的线程级并行,并在硬件调度中动态的调度和执行这些线程。

基于 CUDA 开发的程序代码在实际执行中分为两种,一种是运行在CPU上的宿主代码(Host Code),一种是运行在GPU上的设备代码(Device Code)。不同类型的代码由于其运行的物理位置不同,能够访问到的资源不同,因此对应的运行期组件也分为公共组件、宿主组件和设备组件三个部分,基本上囊括了所有在 GPGPU 开发中所需要的功能和能够使用到的资源接口,开发人员可以通过运行期环境的编程接口实现各种类型的计算。CUDA 所提供的运行期环境是通过驱动来实现各种功能的。
CUDA同时也支持多种语言,但是实际上无论使用那种语言或者接口,指令最终都会被驱动程序转换成PTX(Parallel Thread Execution,并行线程执行,CUDA架构中的指令集,类似于汇编语言)代码,交给GPU核心计算。

在编程的时候一定要注意程序是在那个部分运算的,也就是要注意以下的几个方面:

术语 含义
Host 宿主,指CPU,系统的CPU
Device 设备,指GPU,CPU的协处理器
Block 执行Kernel的一组线程组成一个线程块
线程ID 线程在线程块中的线程号(唯一标识)
Grid 线程块组成的线程网络

CUDA处理流程

  • 从系统内存中复制数据到GPU内存
  • CPU指令驱动GPU运行
  • GPU的每个CUDA核心并行处理
  • GPU将CUDA处理的最终结果返回到系统内存

这里需要注意的是CPU作为主机端只能有一个,GPU作为设备端可以有多个,CPU主要负责逻辑处理,GPU主要负责逻辑密集处理。CUDA包含两个并行逻辑层:block层和thread层,在执行的时候block映射到SM,thread映射到SP(core)

kernel函数的启动

1
gpu_kernel<<<blocks,threads>>>( 参数 );

CUDA编程七步曲

  1. cudaSetDevice(0); //获取设备,一般的机器只有一张GPU卡,或者默认0号GPU的时候可以省略
  2. cudaMalloc((void*) &d_a,sizeof(float)n); //分配显存
  3. cudaMemcpy(d_a,a,sizeof(float)*n,cudaMemoryHostToDevice);
  4. gpu_kernel<<<blocks,threads>>>(***); //kernel函数
  5. cudaMemcpy(a,d_a,sizeof(float)*n,cudaMemcpyDevi
    ceToHost); //数据传输,将GPU的运算结果复制到主内存下
  6. cudaFree(d_a); //释放显存空间
  7. cudaDeviceReset(); //重置设备,可以省略