今天开始重新整理一下CUDA相关内容,重新回头系统学习一下。
大部分参考与谭升大佬的博客:
https://face2ai.com/program-blog/#GPU编程(CUDA)
了解并行
并行计算设计到两个不同的技术领域:
- 计算机架构(硬件)
- 并行程序设计(软件)
硬件主要的目标就是为软件提供更快的计算速度,更低的性能功耗比,硬件结构上支持更快的并行。软件的主要目的是使用当前的硬件压榨出最高的性能,给应用提供更稳定快速的计算结果。
我们传统的计算机结构一般是哈佛体系结构(后来演变出冯·诺依曼结构)主要分成三部分:
- 内存(指令内存,数据内存)
- 中央处理单元(控制单元和算数逻辑单元)
- 输入、输出接口
写并行和串行的最大区别就是,写串行程序可能不需要学习不同的硬件平台,但是写并行程序就需要对硬件有一定的了解了。
写并行程序主要是分解任务,我们一般把一个程序看成是指令和数据的组合,当然并行也可以分为这两种:
- 指令并行
- 数据并行
CUDA非常适合数据并行
数据并行程序设计,第一步就是把数据依据线程进行划分:
不同的数据划分严重影响程序性能,所以针对不同的问题和不同计算机结构,我们要通过和理论和试验共同来决定最终最优的数据划分。
划分不同计算机结构的方法有很多,广泛使用的一种被称为佛林分类法Flynn’s Taxonomy,他根据指令和数据进入CPU的方式分类,分为以下四类:(之前MPI学习过程中也提到过,相通的)
- 单指令单数据SISD(传统串行计算机,386)
- 单指令多数据SIMD(并行架构,比如向量机,所有核心指令唯一,但是数据不同,现在CPU基本都有这类的向量指令)
- 多指令单数据MISD(少见,多个指令围殴一个数据)
- 多指令多数据MIMD(并行架构,多核心,多指令,异步处理多个数据流,从而实现空间上的并行,MIMD多数情况下包含SIMD,就是MIMD有很多计算核,计算核支持SIMD)
为了提高并行的计算能力,我们要从架构上实现下面这些性能提升:
- 降低延迟
- 提高带宽
- 提高吞吐量
延迟是指操作从开始到结束所需要的时间,一般用微秒计算,延迟越低越好。
带宽是单位时间内处理的数据量,一般用MB/s或者GB/s表示。
吞吐量是单位时间内成功处理的运算数量,一般用gflops来表示(十亿次浮点计算),吞吐量和延迟有一定关系,都是反应计算速度的,一个是时间除以运算次数,得到的是单位次数用的时间–延迟,一个是运算次数除以时间,得到的是单位时间执行次数–吞吐量。
CPU与GPU特性:
CPU适合执行复杂的逻辑,比如多分支,其核心比较重(复杂)
GPU适合执行简单的逻辑,大量的数据计算,其吞吐量更高,但是核心比较轻(结构简单)
低并行逻辑复杂的程序适合用CPU
高并行逻辑简单的大数据计算适合GPU
异构计算与CUDA
异构平台,主要是指有不同类型指令集和体系架构计算单元,它可以有CPU,GPU,DSP,ASIC,FPGA等其他处理器构成。
异构架构虽然比传统的同构架构运算量更大,但是其应用复杂度更高,因为要在两个设备上进行计算,控制,传输,这些都需要人为干预,而同构的架构下,硬件部分自己完成控制,不需要人为设计。
CPU可以看做一个指挥者,主机端,host,而完成大量计算的GPU是我们的计算设备,device。
- 左图:一个四核CPU一般有四个ALU,ALU是完成逻辑计算的核心,也是我们平时说四核八核的核,控制单元,缓存也在片上,DRAM是内存,一般不在片上,CPU通过总线访问内存。
- 右图:GPU,绿色小方块是ALU,我们注意红色框内的部分SM,这一组ALU公用一个Control单元和Cache,这个部分相当于一个完整的多核CPU,但是不同的是ALU多了,control部分变小,可见计算能力提升了,控制能力减弱了,所以对于控制(逻辑)复杂的程序,一个GPU的SM是没办法和CPU比较的,但是对了逻辑简单,数据量大的任务,GPU更高效,并且,注意,一个GPU有好多个SM,而且越来越多。
- CPU和GPU之间通过PCIe总线连接,用于传递指令和数据
一个异构应用包含两种以上架构,所以代码也包括不止一部分:
- 主机代码 host
- 设备代码 device
主机代码在主机端运行,被编译成主机架构的机器码,设备端的在设备上执行,被编译成设备架构的机器码,所以主机端的机器码和设备端的机器码是隔离的,自己执行自己的,没办法交换执行。
主机端代码主要是控制设备,完成数据传输等控制类工作,设备端主要的任务就是计算。
因为当没有GPU的时候CPU也能完成这些计算,只是速度会慢很多,所以可以把GPU看成CPU的一个加速设备。
CPU和GPU线程的区别:
CPU线程是重量级实体,操作系统交替执行线程,线程上下文切换花销很大
GPU线程是轻量级的,GPU应用一般包含成千上万的线程,多数在排队状态,线程之间切换基本没有开销。
CPU的核被设计用来尽可能减少一个或两个线程运行时间的延迟,而GPU核则是大量线程,最大幅度提高吞吐量
CUDA:一种异构计算平台
CUDA平台不是单单指软件或者硬件,而是建立在Nvidia GPU上的一整套平台,并扩展出多语言支持:
CUDA nvcc编译器会自动分离你代码里面的不同部分,如图中主机代码用C写成,使用本地的C语言编译器编译,设备端代码,也就是核函数,用CUDA C编写,通过nvcc编译,链接阶段,在内核程序调用或者明显的GPU设备操作时,添加运行时库。
注意:核函数是我们后面主要接触的一段代码,就是设备上执行的程序段
/*
*hello_world.cu
*/
#include<stdio.h>
__global__ void hello_world(void)
{
printf("GPU: Hello world!\n");
}
int main(int argc,char **argv)
{
printf("CPU: Hello world!\n");
hello_world<<<1,10>>>();
cudaDeviceReset();//if no this line ,it can not output hello world from gpu
return 0;
}
- –global-- 是告诉编译器这个是个可以在设备上执行的核函数
- hello_world<<<1,10>>>();这句话C语言中没有’<<<>>>’是对设备进行配置的参数,也是CUDA扩展出来的部分
- cudaDeviceReset();这句话如果没有,则不能正常的运行,因为这句话包含了隐式同步,GPU和CPU执行程序是异步的,核函数调用后成立刻会到主机线程继续,而不管GPU端核函数是否执行完毕,所以上面的程序就是GPU刚开始执行,CPU已经退出程序了,所以我们要等GPU执行完了,再退出主机线程。
一般CUDA程序分成下面这些步骤:
1、分配GPU内存
2、拷贝内存到设备
3、调用CUDA内核函数来执行计算
4、把计算完成数据拷贝回主机端
5、内存销毁