1、GPU架构概述
- GPU是一种众核架构,非常适合解决大规模的并行计算。
- GPU是CPU的协处理器,必须通过PCIe总线与基于CPU的主机(Host)相连来进行操作,形成异构架构,如下图所示。其中CPU为
主机端(Host)
,负责逻辑控制、数据分发,GPU为设备端(Device)
,负责并行数据的密集型计算。其中,ALU为算数运算单元。
- GPU架构是围绕一个流式多处理器
(SM)
的可扩展阵列搭建的。下图是英伟达公司的Fermi架构SM的示意图,SM的关键组件包括GPU核心、共享内存/一级缓存、寄存器文件、加载/存储单元、特殊功能单元和线程束调度器
。 - GPU的每个SM都支持数百个线程并发执行,每个GPU有多个SM,这表示一个GPU可以并发执行数千个线程。
- 当启动一个内核网格
(Grid)
时,它的线程块(block)
被分布在了可用的SM上执行。
- 多个block可能会被分配到同一个SM上。
2、CUDA 并行计算
- CUDA 是英伟达公司推出的通用并行计算平台和编程模型,它利用英伟达的GPU能够实现并行计算。
- CUDA可以通过CUDA加速库、编译器指令、应用编程接口以及标准程序语言的扩展(包括C、C++、Fortan、Python)来使用。
CUDA线程模型
- 线程
(Thread)
是GPU的最小执行单元,能够完成一个逻辑操作,每个线程都有自己的指令地址计数器和寄存器状态,利用自身的数据执行当前的指令。 - 而线程束
(Warp)
是GPU的基本执行单元,包括32个线程,GPU每次调用线程都是以线程束为单位的,在一个线程束中,所有的线程按照单指令多线程(SIMT)
方式执行,即所有线程执行相同的指令。 - 多个线程束位于一个最高维度为3的线程块
(Block)
中,同一个线程块中的所有线程,都可以使用共享内存
来进行通信、同步。 - 线程块又通过一个最高维度为3的网格(Grid)来管理。
CUDA 内存模型
- CUDA内存模型中,有两种类型的存储器:不可编程存储器和可编程存储器,
- 前者并不对开发人员开放存取接口,包括一级缓存和二级缓存;
- 后者可以显式地控制数据在内存空间中的存取,包括寄存器、共享内存、本地内存、常量内存、纹理内存以及全局内存。
- 寄存器:
- 速度最快,分配于每个线程中,数量有限,如果一个核函数使用了超过了限定数量的寄存器,将会溢出到本地内存,降低算法性能。
- 本地内存:
- 本地内存用来存放寄存器溢出的内存,本地内存访问符合高效内存访问要求。
- 共享内存:
- 比本地内存和全局内存有更高的带宽和更低的延迟。它由线程块分配,生命周期伴随着线程块,线程块中的每个线程都可以共享其存储空间。
- 一个块内的线程可以通过共享内存进行通信合作。常用的方式是将全局内存读进共享内存中,而读取的方式是每个线程负责读取某一个位置的数据,读完之后块内的所有线程都能够使用整个共享内存中的数据。
- 读取全局内存到共享内存时要注意同步,在CUDA C中使用线程同步函数
__syncthreads()
来实现同步。 - 核函数中存储在共享内存的变量通过修饰符
__shared__
修饰。 - 常量内存:
- 用修饰符
__constant__
修饰 - 必须在全局空间内和所有核函数之外声明,对同一编译单元的所有线程核函数可见
- 只读
- 纹理内存:
- 只读
- 适合访问二维数据
- 全局内存:
- 是GPU中最大、延迟最高、最常使用的内存,贯穿程序的整个生命周期。
- 对全局内存访问时,必须注意内存访问的两个特性:
对齐内存访问和合并内存访问
。 - 当一个线程束中全部的32个线程访问一个连续的内存块时,满足合并访问,效率非常高。