CUDA C编程(一) - 基于CUDA的异构并行计算

《CUDA C编程权威指南》第一章

并行计算

就我们程序员而言,一个程序包含了指令和数据,对于一个具体问题,我们会尝试将问题进行拆解形成子问题或者子模块,模块之间可能会存在依赖关系,即一个模块的输出会作为另一个模块的输入,这样的关系只能串行。

并行条件

只有相互没有依赖关系的相互独立的模块才能并行执行,所以分析模块间的相关性就显得非常重要。

并行的类型

有两种类型的并行,分别是任务并行和数据并行。任务并行指的是多个可以同时执行的任务或函数;数据并行指的是可以同时处理多个数据,CUDA编程主要处理这一类并行。举个例子,对于不同的相互独立的数据,需要执行相同操作,那么这种情况就属于数据并行,对于我边洗碗边唱歌这就属于任务并行。

数据并行中数据的划分

既然需要让数据的处理并行进行,就需要将不同的数据分配给并行线程,以使其并行。有两种基本的划分方法,分别是块划分和周期划分。以一维数据为例,一个数据块指的是一组连续的数据,通常数据块之间具有相同大小。块划分对每一个数据块分配一个线程,而周期划分会为每一个线程分配多个数据块,因此周期划分的数据块要小于块划分的数据块,周期划分中某一个线程的数据块之间间隔为线程数,从而在数据块属于哪一个线程上形成周期变化。如下图,块的颜色代表了块属于哪一个线程。

通常程序的性能与数据的物理存储以及线程的执行次序、块的大小等有关系。

计算机架构

根据弗林分类法(Flynn’s Taxonomy),分为四种类型。

• SISD:单指令单数据,串行架构,计算机只有一个核心,在每一个时间点都只有一个指令流处理一个数据流。
• SIMD:并行架构,有多个核心,所有核心执行相同指令流,但处理不同的数据流。
• MISD:每个核心使用不同的操作处理相同的数据流,即每个核心通过使用多个指令流处理同一个数据流。
• MIMD:多个核心使用多个指令流异步处理多个数据流。

具体可以参见维基百科。按照内存的组织方式可以分为两类:

• 分布式内存的多节点系统:每个处理器有自己的内存,处理器之间通过网络通信。
• 共享内存的多处理器系统:多个处理器共享同一个物理内存,我们平时使用的笔记本就是这一类型。

GPU是一种拥有很多核心的并行架构,包括多线程、MIMD、SIMD、指令级并行,英伟达公司称这种架构为SIMT(单指令多线程)。GPU可以处理任务并行和数据并行。GPU的核心与CPU的核心差别较大,CPU核心主要用于处理复杂的逻辑控制,而GPU主要偏向于运算和并行。

异构计算

同构计算是指同一架构下的一个或多个处理器来执行一个应用;而异构计算是不同架构下的处理器协同执行一个应用程序,例如CPU处理逻辑控制而GPU负责大规模并行运算。

异构架构

异构架构是由CPU和GPU协同工作而组成的整体。GPU常作为CPU的协处理器,CPU所在的位置被称作主机端,GPU所在的位置被称作设备端,如下图所示

对于我们的并行程序来说,一般包括两个部分,分别是主机代码和设备代码,它们分别在CPU和GPU上运行,在GPU上执行的任务由CPU代码初始化。GPU是与CPU物理上分离的用于加速计算密集型应用的硬件组织,所以GPU也被称为硬件加速器。英伟达公司GPU产品包括了:

• Tegra系列:专为移动端或嵌入式设备设计,Tegra K1包含一个Kepler GPU;
• GeForce系列:面向图形用户;
• Quadro系列:用于专业绘图设计;
• Tesla系列:用于大规模的并行计算,Tesla系列的一些GPU使用了Fermi、Kepler架构,以及目前最先进的Pascal架构。

描述GPU性能的两个重要特征是核心数和显存大小,相应的两个评价GPU性能的指标是峰值计算性能和显存带宽。峰值计算性能定义为每秒能够处理的单精度或双精度浮点运算的数量,单位一般为GFlops(每秒十亿次浮点运算)和TFlops(每秒万亿次浮点运算)。显存带宽是指从显存中读取或写入数据的速度,单位常用GB/s。
英伟达公司使用计算能力(compute capability)去描述整个Tesla系列的GPU加速器的硬件版本,具有相同主版本号(3.x Kepler类架构,2.x Fermi架构)的的设备具有相同的核心架构。计算能力越高,性能越强劲。

异构计算范例

因为CPU和GPU的功能互补性导致了CPU+GPU的异构并行计算架构的发展,一般地在CPU上执行串行部分或任务并行部分,在GPU上执行数据密集型并行部分,为了支持使用CPU+GPU异构系统架构来执行应用程序,NVIDIA设计了一个被称为CUDA的编程模型。

CPU线程和GPU线程

CPU上的线程是重量级的切换缓慢,而GPU线程是轻量级的,CPU的核被设计用来尽可能减少一个或两个线程运行时间的延迟,而GPU的核是用来处理大量并发的、轻量级的线程,以最大限度地提高吞吐量。现在,四核CPU上可以同时运行16个线程,如果CPU支持超线程可支持多至32个线程。普通的NVIDIA GPU在每个多处理器上最多可以并发支持1536个同时活跃的线程。有16个多处理器的GPU,可以并发支持超过24000个同时活跃的线程。

CUDA:一个异构计算平台

CUDA是一种通用的并行计算平台和编程模型,它利用Nvidia GPU并行计算引擎能够很有效地解决数据密集型并行计算应用。CUDA平台可以通过CUDA加速库、编译器指令、应用编程接口以及行业标准程序语言的扩展(包括C、C++、Fortran、Python,如图1-12所示)来使用。基于C和C++的性能,此系列博客以CUDA C编写程序,我们能够通过CUDA提供的API管理设备、显存和其他任务,而且CUDA还是一个可扩展的编程模型,它使程序能对有不同数量核的GPU明显地扩展其并行性。
CUDA提供了两层API来管理GPU设备和组织线程,如下图所示

其中CUDA驱动API是底层API较难编程,但能提供更多的功能,CUDA运行时API是基于驱动API开发的,是较高级的API,这两种API是相互排斥的,你只能在你的程序中使用其中一个API,后面都使用运行时API。

编译

NVIDIA的CUDA nvcc编译器在编译过程中会将设备代码从主机代码中分离出来。主机代码是标准的C代码,使用C编译器进行编译。设备代码,也就是核函数,是用扩展的CUDA C语言编写的,设备代码通过nvcc进行编译。在链接阶段,在内核程序调用和显示GPU设备操作中添加CUDA运行时库,如下图所示

CUDA可以自己拓展编程语言,这样的生态是靠大家来创建啊,CUDA nvcc编译器是以广泛使用LLVM开源编译系统为基础的,在GPU加速器的支持下,通过使用CUDA编译器SDK,我们可以创建或扩展编程语言。

第一行GPU代码

注意使用nvcc编译器对于C/C++代码部分会调用gcc编译器,可能存在gcc版本问题,根据错误提示,在/usr/bin目录下更改gcc链接就行了。

#include <stdio.h>
//__global__告诉编译器这个函数将会从CPU上调用,在GPU上执行
__global__ void helloFromGPU(void)
{
    printf("Hello world from GPU!\n");
}
int main(void)
{
    printf("Hello world from CPU!\n");
    //三重尖括号意味着从主线程到设备端代码的调用
    //第一个参数表示启用多少个线程块,第二个参数表示每个线程块中启用多少个线程
    //所有线程执行相同的代码
    helloFromGPU <<<1,10>>>();
    //显式地清空进程中与当前设备相关的所有资源
    cudaDeviceReset();
    return 0;
}

在编译时需要指出的是,因为不同显卡的架构不同,编译的指令也就不同了,你需要确定你的显卡的计算能力,如果为7,那么需要选项-arch sm_70,我的显卡是TITAN V,所以编译命令为

nvcc -arch sm_70 HelloWorld.cu -o HelloWorld

这里得到了可执行文件,直接使用以下命令就可运行

./HelloWorld

最后输出

CUDA编程结构

一个典型的CUDA编程结构包括五个主要步骤:

• 分配GPU显存
• 从内存中拷贝数据到GPU显存上
• 执行内核函数
• 从GPU显存拷贝数据回内存中
• 释放GPU显存

数据局部性

数据局部性指的是刚使用过的数据和空间上较接近的数据容易被重用,现代的CPU使用高速缓存对这一部分数据进行优化,以减少访问内存的延迟。在GPU中,我们需要进行优化,在CUDA编程模型中使用的共享内存(一个特殊的内存),共享内存可以视为一个被软件管理的高速缓存,通过为主内存节省带宽来大幅度提高运行速度。有了共享内存,我们可以直接控制代码的数据局部性。
当用CUDA C编写程序时,实际上你只编写了被单个线程调用的一小段串行代码。GPU处理这个内核函数,然后通过启动成千上万个线程来实现并行化,所有的线程都执行相同的计算(数据不同)。因为CUDA C是C语言的扩展,通常可以直接将C程序移植到CUDA C程序中。概念上,因为需要并行的计算,往往是循环体,所以剥离代码中的循环后产生CUDA C实现的内核代码(内核代码被多线程并行执行)。

CUDA开发环境

CUDA核中有3个关键抽象:线程组的层次结构,内存的层次结构以及障碍同步。CUDA终究也必须是一个底层的平台,才能拥有更强的性能,所以我们的学习目标应是GPU架构的基础及掌握CUDA开发工具和环境。CUDA C开发环境提供了很多工具供程序员使用,包括:

• NVIDIA Nsight集成开发环境
• CUDA-GDB命令行调试器
• 用于性能分析的可视化和命令行分析器
• CUDA-MEMCHECK内存分析器
• GPU设备管理工具

总结

CPU+GPU的异构系统在高性能计算领域已经成为主流。这种变化使并行设计范例有了根本性转变:在GPU上执行数据并行工作,而在CPU上执行串行和任务并行工作。

习题

2. 从hello.cu中移除cudaDeviceReset函数,然后编译运行,看看会发生什么。
   运行结果是:
   
3. 用cudaDeviceSynchronize函数来替换hello.cu中的cudaDeviceReset函数,然后编译运行,看看会发生什么。
   
4. 参考1.3节,从编译器命令行中移除设备架构标志,然后按照下面的方式进行编译,看看会发生什么。

nvcc hello.cu -o hello

          会提示错误

6. 为执行核函数的每个线程提供了一个唯一的线程ID,通过内置变量threadIdx.x可以在内核中对线程进行访问。在hello.cu中修改核函数的线程索引,使输出如下:
   
   执行结果如下图所示
   
   代码:

#include <stdio.h>
//__global__告诉编译器这个函数将会从CPU上调用,在GPU上执行
__global__ void helloFromGPU(void)
{
    //threadIdx.x为整型数
    printf("Hello world from GPU thread %d!\n",threadIdx.x);
}
int main(void)
{
    printf("Hello world from CPU!\n");
    //三重尖括号意味着从主线程到设备端代码的调用
    //第一个参数表示启用多少个线程块,第二个参数表示每个线程块中启用多少个线程
    //所有线程执行相同的代码
    helloFromGPU <<<1,10>>>();
    //显式地清空进程中与当前设备相关的所有资源
    //cudaDeviceReset();
    cudaDeviceSynchronize();
    return 0;
}