深度学习利器-gpu介绍
近代史科技发展日新月异，摩尔定律从中显威，各种底层技术层出不穷，但是纵观科技发展史，几乎所有的新兴学科的发展背后都有一个字——“钱”！
作为近年来最火热的行业——人工智能，在烧钱方面同样不遑多让。众所周知，人工智能的训练和推理都需要海量的高性能计算，做深度学习的朋友都知道，现今深度学习领域的sota模型往往需要巨大的显存空间，这直接导致了深度学习的研究者们需要配置更强劲的 gpu 设备，否则就会分分钟面临显存与算力不足的窘境。大型企业或者研究机构可以一掷千金部署 hpc，比如近年来特别热点的gpt3等sota模型，可以这么说，想要跑下来、跑的快可能没有个上千万都不要去想。
当今的 ai 模型面临着对话式 ai 等更高层次的挑战，这促使其复杂度呈爆炸式增长。训练这些模型需要大规模的计算能力和可扩展性。
以英伟达目前的最先进的a100为例，nvidia a100 的 tensor core 借助 tensor 浮点运算 (tf32) 精度，可提供比上一代 nvidia volta 高 20 倍之多的性能，并且无需更改代码；若使用自动混合精度和 fp16，性能可进一步提升 2 倍。
2048 个 a100 gpu 可在一分钟内大规模处理 bert 之类的训练工作负载，这是训练时间的世界纪录。
对于具有庞大数据表的超大型模型（例如用于推荐系统的 dlrm），a100 80gb 可为每个节点提供高达 1.3 tb 的统一显存，而且速度比 a100 40gb 快高达 3 倍。
nvidia 产品的领先地位在 mlperf 这个行业级 ai 训练基准测试中得到印证，创下多项性能纪录。
a100 80gb相对于a100 40g来说在gpu芯片上没变化，依然是a100核心，6912个cuda核心，加速频率1.41ghz，fp32性能19.5tflops，fp64性能9.7tflops，int8性能624tops，tdp 400w。变化的主要是显存，之前是40gb，hbm2规格的，带宽1.6tb/s，现在升级到了80gb，显存类型也变成了更先进的hbm2e，频率从2.4gbps提升到3.2gbps，使得带宽从1.6tb/s提升到2tb/s。
对于大众来说，cpu是用来处理计算的，而gpu是用来处理图像渲染的，从能力范围来说相对来说cpu可能更加强大，那么深度学习为什么可以使用gpu来进行加速，性能获得巨大的提升呢，下面和大家娓娓道来，可以先泡一杯茶，慢慢来看。
综合来说，大家可以这么理解，举个职场的例子，相信很多人都遇到过，cpu是一个十八班武义精通汉子，但是在职场里面由于其精通的太多，所以导致什么事情都分到了他那里，所以他在疯狂的流转，大家没注意到他做到的，只注意到他没做到的。但是gpu呢，他比较矫情，也有一定的能力，但是能力范围有限，而且他挑活，相对于cpu的任劳任怨，他只想做能出成绩的，对于其他一概不理。延伸来说吧，还有一些人是sbpu，啥叫sbpu呢，就是能力也不咋地，抢别人的功劳一流，但是自己屁也做不出来，靠着不光彩的手段和龌龊的计量，占据别人的功劳为己用，然后狂吹ppt。额，跑题了，赶紧回到正题。
2.1 cpu介绍
中央处理器（cpu），是电子计算机的主要设备之一，电脑中的核心配件。其功能主要是解释计算机指令以及处理计算机软件中的数据。cpu是计算机中负责读取指令，对指令译码并执行指令的核心部件。中央处理器主要包括两个部分，即控制器、运算器，其中还包括高速缓冲存储器及实现它们之间联系的数据、控制的总线。电子计算机三大核心部件就是cpu、内部存储器、输入/输出设备。中央处理器的功效主要为处理指令、执行操作、控制时间、处理数据。
在计算机体系结构中，cpu 是对计算机的所有硬件资源（如存储器、输入输出单元） 进行控制调配、执行通用运算的核心硬件单元。cpu 是计算机的运算和控制核心。计算机系统中所有软件层的操作，最终都将通过指令集映射为cpu的操作。
本质：cpu 由专为顺序串行处理而优化的几个核心组成，当然也可以通过多核与超线程等技术进行并行，但是受限于cpu的复杂工作流程、功耗、散热、尺寸等方面的影响，整体cpu的核心不会太多，这点也是gpu的优势。
2.2 gpu介绍
图形处理器（英语：graphics processing unit，缩写：gpu），又称显示核心、视觉处理器、显示芯片，是一种专门在个人电脑、工作站、游戏机和一些移动设备（如平板电脑、智能手机等）上做图像和图形相关运算工作的微处理器。 [1]
gpu使显卡减少了对cpu的依赖，并进行部分原本cpu的工作，尤其是在3d图形处理时gpu所采用的核心技术有硬件t&l（几何转换和光照处理）、立方环境材质贴图和顶点混合、纹理压缩和凹凸映射贴图、双重纹理四像素256位渲染引擎等，而硬件t&l技术可以说是gpu的标志。gpu的生产商主要有nvidia和ati。
2.3 cpu与gpu的区别
首先，针对gpu 和 cpu的各自特点进行分析，阐述下gpu相对于cpu为何在并行计算上具备优势。
- 任务模式
- cpu 由专为顺序串行处理而优化的几个核心组成
- gpu 则拥有一个由数以千计的更小、更高效的核心（专为同时处理多重任务而设计）组成的大规模并行计算架构。同时cpu相当的一部分时间在执行外设的中断、进程的切换等任务，而gpu有更多的时间并行计算。
- 功能定位
- gpu在渲染画面时需要同时渲染数以百万记的顶点或三角形，故gpu的设计是可以充分支持并行计算。
- cpu不但要承担计算任务还有承担逻辑控制等任务。
- 系统集成
- gpu作为一类外插设备，在尺寸、功率、散热、兼容性等方面的限制远远小于cpu，这样可以让gpu有较大的显存和带宽。
2.4 gpu在深度学习大放异彩的关键 – cuda
gpu已经存在了很长时间，那么为何最近几年才开始大放异彩呢。个人认为主要原因有以下两个：
- 最近几年神经网络深度学习对算力的要求不断攀升，特别是cv、nlp领域的大量超算操作。
- 以前nvidia等厂商没有提供gpu – 深度学习之间的桥梁，把gpu真正用到深度虚线上有很大的成本，直至cuda的出现，顺利的打通了这个桥梁。给gpu的盛行，提供了坚实的基座。
在软件层，通过cuda抽象成了统一的编程接口并提供c/c /python/java等多种编程语言的支持。cuda的这一层抽象非常重要，因为有了这层抽象，向下开发者的代码可在不同的硬件平台上快速迁移；向上，在cuda基础上封装了用于科学计算的cublas（兼容blas接口），用于深度学习的cudnn等中间件和代码库。这些中间件对于tensorflow，pytorch这一类的深度学习框架非常重要，可以更容易地使用cuda和底层硬件进行机器学习的计算任务。
3.1 gpu的物理体系结构
仍然以英伟达的a100 gpu为例，从整体上来说a100具备6192个core，所以大家可以这么理解，他是一个超级的cpu，可以并行执行6192个核心。如下图
gpu的基础单位是sm，实际上在 nvidia 的 gpu 里，最基本的处理单元是所谓的 sp(streaming processor)，而一颗 nvidia 的 gpu 里，会有非常多的 sp 可以同时做计算；而数个 sp 会在附加一些其他单元，一起组成一个 sm(streaming multiprocessor)。几个 sm 则会在组成所谓的 tpc(texture processing clusters)。
这里我们暂时把sm可以看成是gpu计算调度的一个基本单位。一个sm具备64个用于计算的core。针对a100来说，整个gpu上面有个108条好汉，额，是108个sm，所以整体的a100的core的数量就是 64 * 108 = 6192。看看，够强劲吧，1个卡可以同时跑6192个并行的计算任务。
针对这个sm的架构图，再简单进行下补充：
- 最上面是pcie层，通过pcie接口以外设的方式集成到服务器上。
- 绿色的部分是gpu的计算核心，比如a100具备6192个core。
- 中间蓝色部分是l2缓存
- nvlink是多个gpu间进行通信的组件，会对gpu之间的通信做些优化，减轻cpu负担，提升传输效率，这点在分布式训练中会用的比较多。
- 两侧的hbm2就是显存，目前的a100的显存有两种40g and 80g，有钱就买80g吧，差一半的显存呢
既然sm这么重要，那么我们就看看sm的硬件。
sm的相关结构：
3.2 gpu的逻辑体系结构
如果把 cuda 的 grid - block - thread 架构对应到实际的硬件上的话，会类似对应成 gpu - streaming multiprocessor - streaming processor;一整个 grid 会直接丢给 gpu 来执行，而 block 大致就是对应到 sm，thread 则大致对应到 sp。当然，这个讲法并不是很精确，只是一个简单的比喻而已。
- kernel：thread执行的内容/代码/函数
- thread：执行kernel的最小单元，调度到cuda core中执行
- warp：gpu 有很多 streaming multiprocessors, 用来管理调度 thread block
- thread block：。多个thread组合，被调度到sm中执行。一个sm可以同时执行多个thread block，但是一个thread block只能被调度到一个sm上。gpu 有很多 streaming multiprocessors, 用来管理调度 thread block。
- grid：多个thread block的组合，被调度到整个gpu中执行
同时，thread、thread block和grid由于所处层次不同，他们可以访问的存储资源也不同。如thread只能访问自身的寄存器，thread block可以访问sm中的l1缓存，而grid则可以访问l2缓存和更大的hbm显存。
gpu的运行主要是四步：
- 从host拷贝数据到device
- cpu发送数据处理执行给gpu；
- 把需要device执行的kernel函数发射给device
- 从device拷贝计算结果到host
4.1 cuda编程实战
// device code
__global__ void vecadd(float* a, float* b, float* c, int n)
{
int i = blockdim.x * blockidx.x threadidx.x;
if (i < n)
c[i] = a[i] b[i];
}
// host code
int main()
{
int n = ...;
size_t size = n * sizeof(float);
// allocate input vectors h_a and h_b in host memory
float* h_a = (float*)malloc(size);
float* h_b = (float*)malloc(size);
float* h_c = (float*)malloc(size);
// initialize input vectors
...
// allocate vectors in device memory
float* d_a;
cudamalloc(&d_a, size);
float* d_b;
cudamalloc(&d_b, size);
float* d_c;
cudamalloc(&d_c, size);
// copy vectors from host memory to device memory
cudamemcpy(d_a, h_a, size, cudamemcpyhosttodevice);
cudamemcpy(d_b, h_b, size, cudamemcpyhosttodevice);
// invoke kernel
int threadsperblock = 256;
int blockspergrid =
(n threadsperblock - 1) / threadsperblock;
vecadd<<>>(d_a, d_b, d_c, n);
// copy result from device memory to host memory
// h_c contains the result in host memory
cudamemcpy(h_c, d_c, size, cudamemcpydevicetohost);
// free device memory
cudafree(d_a);
cudafree(d_b);
cudafree(d_c);
// free host memory
...
}
这里是如何体现出gpu的威力呢，举个例子：如果gpu中的每个核心都有唯一的id，比如0号核心执行c[0] = a[0] b[0]，127号核心执行c[127] = a[127] b[127]就好了。
int i = blockdim.x * blockidx.x threadidx.x;主要是计算线程id，那么关于线程的编码是如何呢，请看下面的简介。
那么关于线程的编码是如何呢，请看下面的简介。
- threadidx是一个uint3类型，表示一个线程的索引。
- blockidx是一个uint3类型，表示一个线程块的索引，一个线程块中通常有多个线程。
- blockdim是一个dim3类型，表示线程块的大小。
- griddim是一个dim3类型，表示网格的大小，一个网格中通常有多个线程块。
- 下面这张图比较清晰的表示的几个概念的关系：
4.2 基于gpu的加速
整个gpu的东西还是比较多的，这块我就先简单写下吧，后续会在开个专题，这里先简单列几条：
- 算子融合
- 空域滤波
- high level graph opt
- cpu & gpu filter