异构计算–CUDA架构

大家好，欢迎来到IT知识分享网。

1.CUDA是什么？

  CUDA(Compute Unified Device Architecture)，是显卡厂商NVIDIA推出的运算平台，是一种通用并行计算架构，该架构使GPU能够解决复杂的计算问题。说白了就是我们可以使用GPU来并行完成像神经网络、图像处理算法这些在CPU上跑起来比较吃力的程序。通过GPU和高并行，我们可以大大提高这些算法的运行速度。

2.CPU&CUDA架构

  处理器结构有2个指标要经常考虑的：延迟和吞吐。延迟指从发出指令到返回最终结果中间经历的时间间隔；吞吐指单位时间内处理的指令的条数。由于CPU以处理计算和控制为主要任务，所以设计理念是延迟导向内核；由于GPU以并行处理为主要任务，所以设计理念是吞吐导向内核。

（1）CPU

  CPU遵循的是冯·诺依曼架构，其核心是存储程序/数据、串行顺序执行。因此CPU的架构中需要大量的空间去放置存储单元（Cache）和控制单元（Control），相比之下计算单元（ALU）只占据了很小的一部分，所以CPU在进行大规模并行计算方面受到限制，相对而言更擅长于处理逻辑控制。CPU无法做到大量数据并行计算的能力，但GPU可以。

（2）GPU

3.NVIDIA显卡硬件架构：SM、SP、Warp

流处理器SP（streaming processor）
最基本的处理单元，也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。
SM（streaming multiprocessor）
多个SP加上其他的一些资源组成一个SM，也叫GPU大核，其他资源如：warp scheduler，register，shared memory等。SM可以看做GPU的心脏（对比CPU核心），register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。如下图是一个SM的基本组成，其中每个绿色小块代表一个SP。

每个SM包含的SP数量依据GPU架构而不同，Fermi架构GF100是32个，GF10X是48个，Kepler架构都是192个，Maxwell都是128个。当一个kernel启动后，thread会被分配到很多SM中执行。大量的thread可能会被分配到不同的SM，但是同一个block中的thread必然在同一个SM中并行执行。
Warp调度
一个SP可以执行一个thread，但是实际上并不是所有的thread能够在同一时刻执行。Nvidia把32个threads组成一个warp，warp是调度和运行的基本单元。warp中所有threads并行的执行相同的指令。warp由SM的硬件warp scheduler负责调度，一个SM同一个时刻可以执行多个warp，这取决于warp scheduler的数量。目前每个warp包含32个threads（Nvidia保留修改数量的权利）。

4.软件架构：Kernel、Grid、Block

具体到我们如何调用GPU上的线程实现我们的算法，则是通过Kernel实现的。在GPU上调用的函数成为CUDA核函数（Kernel function），核函数会被GPU上的多个线程执行。我们可以通过如下方式来定义一个kernel：

kernel_function<<<grid, block>>>(param1, param2, param3....); 

Block： 一个grid由许多block组成，block由许多线程组成，同样可以有一维、二维或者三维。block内部的多个线程可以同步（synchronize），可访问共享内存（share memory）。

5.CUDA内存模型

CUDA中的内存模型分为以下几个层次：

• 每个线程都用自己的registers（寄存器）
• 每个线程都有自己的local memory（局部内存）
• 每个线程块内都有自己的shared memory（共享内存），所有线程块内的所有线程共享这段内存资源
• 每个grid都有自己的global memory（全局内存），不同线程块的线程都可使用
• 每个grid都有自己的constant memory（常量内存）和texture memory（纹理内存），），不同线程块的线程都可使用
• 线程访问这几类存储器的速度是register > local memory >shared memory > global memory

下面这幅图表示就是这些内存在计算机架构中的所在层次。

6.异构计算

7.OpenCL与CUDA的关系