【CUDA 基础】1.1 异构计算与CUDA

发表于 | 分类于 ， | 评论数： | 阅读次数：

Abstract: 介绍异构计算和CUDA概述，完成GPU输出Hello world！
Keywords: 异构计算，CUDA

异构计算与CUDA

异构计算

异构计算，首先必须了解什么是异构，不同的计算机架构就是异构，上文书我们讲过计算机架构了，就是为了引出异构的概念，按照指令集划分或者按照内存结构划分，但是我觉得只要两片CPU型号不一样就应该叫异构（这个想法先保留，对错不确定）。
GPU本来的任务是做图形图像的，也就是把数据处理成图形图像，图像有个特点就是并行度很高，基本上一定距离意外的像素点之间的计算是独立的，所以属于并行任务。
GPU之前是不可编程的，或者说不对用户开放的，人家本来是做图形计算控制显示器的，虽然对用户不可编程，但是你只要把硬件卖给了我，就由不得你了，然后就有hacker开始想办法给GPU编程，来帮助他们完成规模较大的运算，于是他们研究着色语言或者图形处理原语来和GPU对话。后来黄老板发现了这个是个新的功能啊，然后就让人开发了一套平台，CUDA，然后深度学习火了，顺带着，CUDA也火到爆炸。
刚刚最新消息，英伟达新版本GPU架构会被命名为Turing，一丝欣慰，发自内心深处地敬那些为世界进步做出了杰出贡献的人们，他们是人类未来的希望。

x86 CPU+GPU的这种异构应该是最常见的，也有CPU+FPGA，CPU+DSP等各种各样的组合，CPU+GPU在每个笔记本或者台式机上都能找到。当然超级计算机大部分也采用异构计算的方式来提高吞吐量。
异构架构虽然比传统的同构架构运算量更大，但是其应用复杂度更高，因为要在两个设备上进行计算，控制，传输，这些都需要人为干预，而同构的架构下，硬件部分自己完成控制，不需要人为设计。

异构架构

举一个我用的工作站的构成，我使用的是一台 intel i7-4790 CPU加上两台Titan x GPU构成的工作站，GPU插在主板的PCIe卡口上，运行程序的时候，CPU像是一个控制者，指挥两台Titan完成工作后进行汇总，和下一步工作安排，所以CPU我们可以把它看做一个指挥者，主机端，host，而完成大量计算的GPU是我们的计算设备，device。

上面这张图能大致反应CPU和GPU的架构不同。

  • 左图：一个四核CPU一般有四个ALU，ALU是完成逻辑计算的核心，也是我们平时说四核八核的核，控制单元，缓存也在片上，DRAM是内存，一般不在片上，CPU通过总线访问内存。
  • 右图：GPU，绿色小方块是ALU，我们注意红色框内的部分SM，这一组ALU公用一个Control单元和Cache，这个部分相当于一个完整的多核CPU，但是不同的是ALU多了，control部分变小，可见计算能力提升了，控制能力减弱了，所以对于控制（逻辑）复杂的程序，一个GPU的SM是没办法和CPU比较的，但是对了逻辑简单，数据量大的任务，GPU更搞笑，并且，注意，一个GPU有好多个SM，而且越来越多。

CPU和GPU之间通过PCIe总线连接，用于传递指令和数据，这部分也是后面要讨论的性能瓶颈之一。
一个异构应用包含两种以上架构，所以代码也包括不止一部分：

  • 主机代码
  • 设备代码

主机代码在主机端运行，被编译成主机架构的机器码，设备端的在设备上执行，被编译成设备架构的机器码，所以主机端的机器码和设备端的机器码是隔离的，自己执行自己的，没办法交换执行。
主机端代码主要是控制设备，完成数据传输等控制类工作，设备端主要的任务就是计算。
因为当没有GPU的时候CPU也能完成这些计算，只是速度会慢很多，所以可以把GPU看成CPU的一个加速设备。
NVIDIA目前的计算平台（不是架构）有：

  • Tegra
  • Geforce
  • Quadro
  • Tesla

每个平太针对不同的应用场景，比如Tegra用于嵌入式，Geforce是我们平时打游戏用到，Tesla是我们昨天租的那台腾讯云的，主要用于计算。

上面是根据应用场景分类的几种平台。

衡量GPU计算能力的主要靠下面两种容量特征：

  • CUDA核心数量（越多越好）
  • 内存大小（越大越好）

相应的也有计算能力的性能指标:

  • 峰值计算能力
  • 内存带宽

nvidia自己有一套描述GPU计算能力的代码，其名字就是“计算能力”，主要区分不同的架构，早其架构的计算能力不一定比新架构的计算能力强

计算能力 架构名
1.x Tesla
2.x Fermi
3.x Kepler
4.x Maxwell
5.x Pascal
6.x Volta

这里的Tesla架构，与上面的Tesla平台不同，不要混淆，一个是平台名字，一个是架构名字

范例

CPU和GPU相互配合，各有所长，各有所短，不能说GPU就是比CPU强这种幼稚的话：

低并行逻辑复杂的程序适合用CPU
高并行逻辑简单的大数据计算适合GPU

一个程序可以进行如下分解，串行部分和并行不分：

CPU和GPU线程的区别：

  1. CPU线程是重量级实体，操作系统交替执行线程，线程上下文切换花销很大
  2. GPU线程是轻量级的，GPU应用一般包含成千上万的线程，多数在排队状态，线程之间切换基本没有开销。
  3. CPU的核被设计用来尽可能减少一个或两个线程运行时间的延迟，而GPU核则是大量线程，最大幅度提高吞吐量

CUDA：一种异构计算平台

CUDA平台不是单单指软件或者硬件，而是建立在Nvidia GPU上的一整套平台，并扩展出多语言支持

CUDA C 是标准ANSI C语言的扩展，扩展出一些语法和关键字来编写设备端代码，而且CUDA库本身提供了大量API来操作设备完成计算。

对于API也有两种不同的层次，一种相对交高层，一种相对底层。

  • CUDA驱动API
  • CUDA运行时API

驱动API是低级的API，使用相对困难，运行时API是高级API使用简单，其实现基于驱动API。
这两种API是互斥的，也就是你只能用一个，两者之间的函数不可以混合调用，只能用其中的一个库。

一个CUDA应用通常可以分解为两部分，

  • CPU 主机端代码
  • GPU 设备端代码

CUDA nvcc编译器会自动分离你代码里面的不同部分，如图中主机代码用C写成，使用本地的C语言编译器编译，设备端代码，也就是核函数，用CUDA C编写，通过nvcc编译，链接阶段，在内核程序调用或者明显的GPU设备操作时，添加运行时库。

注意：核函数是我们后面主要接触的一段代码，就是设备上执行的程序段

nvcc 是从LLVM开源编译系统为基础开发的。

CUDA工具箱提供编译器，数学库，调试优化等工具，当然CUDA的文档是相当完善的，可以去查阅，当然在我们基本了解基础结构的情况下，直接上来看文档会变得机械化。

“Hello World!”

Hello World是所有程序初学者都非常喜欢的，之前GPU是不能printf的，我当时就很懵，GPU是个做显示的设备，为啥不能输出，后来就可以直接在CUDA核里面打印信息了，我们写下面程序

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
/*
*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;
}

完整代码可以从以下项目中clone：

https://github.com/Tony-Tan/CUDA_Freshman

简单介绍其中几个关键字

1
__global__

是告诉编译器这个是个可以在设备上执行的核函数

1
hello_world<<<1,10>>>();

这句话C语言中没有’<<<>>>’是对设备进行配置的参数，也是CUDA扩展出来的部分。

1
cudaDeviceReset();

这句话如果没有，则不能正常的运行，因为这句话包含了隐式同步，GPU和CPU执行程序是异步的，核函数调用后成立刻会到主机线程继续，而不管GPU端核函数是否执行完毕，所以上面的程序就是GPU刚开始执行，CPU已经退出程序了，所以我们要等GPU执行完了，再退出主机线程。
一般CUDA程序分成下面这些步骤：

  1. 分配GPU内存
  2. 拷贝内存到设备
  3. 调用CUDA内核函数来执行计算
  4. 把计算完成数据拷贝回主机端
  5. 内存销毁

上面的hello world只到第三步，没有内存交换。

CUDA C难么？

CPU与GPU的编程主要区别在于对GPU架构的熟悉程度，理解机器的结构是对编程效率影响非常大的一部分，了解你的机器，才能写出更优美的代码，而目前计算设备的架构决定了局部性将会严重影响效率。
数据局部性分两种

  • 空间局部性
  • 时间局部性

这个两个性质告诉我们，当一个数据被使用，其附近的数据将会很快被使用，当一个数据刚被使用，则随着时间继续其被再次使用的可能性降低，数据可能被重复使用。

CUDA中有两个模型是决定性能的：

  • 内存层次结构
  • 线程层次结构

CUDA C写核函数的时候我们只写一小段串行代码，但是这段代码被成千上万的线程执行，所有线程执行的代码都是相同的，CUDA编程模型提供了一个层次化的组织线程，直接影响GPU上的执行顺序。

CUDA抽象了硬件实现：

  1. 线程组的层次结构
  2. 内存的层次结构
  3. 障碍同步

这些都是我们后面要研究的，线程，内存是主要研究的对象，我们能用到的工具相当丰富，NVIDIA为我们提供了：

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

总结

本文从总体上粗略的介绍了CUDA这种高效的异构计算平台，并且概括了我们的将要遇到的苦难和使用到的工具，当我们学会的CUDA，那么编写高效异构计算就会像我们写串行程序一样流畅。
祝大家新年快乐。

相关文章
0%