【CUDA编程】学习笔记（三） CUDA软件抽象机制

一、多线程

（1）设计理念

CPU是面向延迟的核心，尽可能快完成计算任务
GPU是面向吞吐量的核心，通过多个线程实现在某一段时间内的高吞吐率
打个比喻，前者是一架飞机，后者是一千辆牛车

可以看到，CPU的Cache更大，通过复杂的控制逻辑实现性能优化，而GPU的控制逻辑相对简单，通过多线程来实现高吞吐率

（2）多线程

CPU中的SIMD
✓所有内核同时执行相同的指令，但数据不同
✓类似于CRAY超级计算机上的矢量计算
✓例如：SSE4，AVX指令集
SMX中的SIMT
✓多线程CUDA核心
✓每个SMX上的线程执行组共享相同指令
✓细粒度并行
✓自然用于图形处理和大量科学计算
✓SIMT也是多核芯片的自然选择简化每个核心
线程：具有自己的PC和数据的指令流
✓拥有专用寄存器，专用内存，程序计数器和线程执行状态
✓线程级并行（TLP）：利用线程之间固有的并行性
多线程
✓多线程共享功能单元1 处理器通过重叠
✓处理器必须复制每个线程的独立状态，例如，寄存器文件的单独副本，单独的PC
✓通常，用于快速线程切换的硬件
✓要共享的内存
✓解决内存访问停顿

（3）GPU多线程

GPU是一种多线程处理器，例如GTX8800支持96个线程
取值→译码→执行→内存→回写
显然，内存读写是最耗时的阶段，在这个停滞时期进行其他线程的计算可以提高性能

二、CUDA抽象

CPU诞生于2007年，是一种并行计算平台和编程模型，可通过利用GPU的强大功能显着提高计算性能

CUDA作为并行计算平台
✓语言：基于C的CUDA C，带有一些扩展（广泛的C ++支持）
✓编辑器：Eclipse / VisualStudio
✓编译器：nvcc
✓SDK：CUDA工具包，库，样本
✓Profiler＆Debugger：Nsight
CUDA作为编程模型
✓GPU硬件的软件抽象
✓独立于操作系统，CPU，Nvidia GPU

CUDA虚拟化物理硬件
✓线程是虚拟化的CUDA核心（寄存器，PC，状态）
✓block是虚拟化流式多处理器（SM）（线程，共享内存）
无需抢占计划到物理硬件上
✓线程/块启动并运行到完成/暂停
✓块应该是独立的

（1）并发线程的层次结构

第一个参数是block数量，第二个参数是每个block的线程数

线程对应于CUDA核，block对应于SM，而Grid对应于多个SM

（2）协作线程的共享内存模型

每个线程有自己的私有内存和寄存器，同一个block内有共享内存，grid间有全局内存

（3）轻量级同步

全局同步

完成一个kernel再进行下一个
线程同步

kernel内部完成一个线程再进行下一个
竞争
原子操作
atomic{Add, Sub, Exch, Min, Max, Inc, Dec, CAS, And, Or, Xor}

原子操作可以避免竞争，但是执行速度略慢

三、Kernel执行

（1）Kernel函数

使用__global__修饰，返回值为void
__device__和__host__可以一起使用

KernelFunc<<<gridDim, blockDim, nSMem,iStream>>>(args);
gridDim是网格的维度
blockDim是每个块的维度
其中这些参数可以是dim3类型，要注意下面这些维度限制

（2）线程索引

对于一维情况：
blockIdx.x * blockDim.x + threadIdx.x

对于二维情况：
i = blockIdx.x * blockDim.x + threadIdx.x
j = blockIdx.y * blockDim.y + threadIdx.y
idx = j * gridDim.x * blockDim.x + i

如果数据无法恰好比配，则有上面的情况。
对于三位情况图示如下，推导类似于二维情况。

（3）设备占用率

✓每个多处理器的活动warp数与同时在多处理器上可以激活的最大warp数之比
✓分母是一个常数，仅取决于设备的计算能力
✓该表达式的分子是一个函数与以下内容有关：
•计算能力
•每个块的线程数
•每个线程的寄存器数
•每个块的共享内存和共享内存配置

四、线程束调度

在较低级别，在GPU内：
✓执行内核的每个块在SMX上执行，一个SM最大由16个block共享
✓如果块的数量超过SMX的数量，那么如果有足够的寄存器，则每个SMX上将一次运行多个块和共享内存，其他人将在队列中等待并稍后执行
✓一个块中的所有线程都可以访问本地共享内存，但无法看到其他块正在做什么（即使它们在同一个SMX上）
✓不保证块执行的顺序

block映射到物理GPU上按照每32个线程一组，即一个线程束，几乎是同时运行同时结束，可以看到warp0处在执行状态

当进行访存时warp0处于挂起状态，同时warp1进入执行状态，此时warp1如果也需要访存也将挂起

当4个warp都挂起时，硬件此时会空闲

当某个warp访存完成后，则会进入执行阶段

五、CUDA开发工具和生态圈

（1）软件架构

从图中可以发现，CUDA应用有多层结构，而CUDA编程本身其实是基于其中“CUDA Runtime”层次。

（2）编译过程

从图中可以看到，编写完.cu文件后，nvcc编译成GPU与CPU运行的代码，然后合并，经过Host编译器生成Host执行文件，然后运行，并通过CUDA驱动。运行涉及到CUDA的一些库，cudart.lib和cuda.lib。

（3）NVCC

常用命令行选项：
-g：产生可调式代码，调试模式下必须的
-G：产生可调式的设备代码
-lineinfo：为设备代码产生行number信息
-o outfile：指定输出文件的位置和名称
-include xxx.h：指定预处理和编译时预先需要包含的头文件
-l yyy.lib：指定链接时需要的库文件
-arch sm_??：指定GPU架构

（4）CUDA软件生态

数值分析工具：Matlab, ArrayFire, Mathematica, LabView, Jacket
GPU加速库：Thrust, cuDNN, cuFFT, cuBLAS
语言和接口：CUDA Toolkit, OpenACC, CUDA Fortran, PyCUDA, PGI compilers
性能分析工具：Nsight, NVIDIA Visual Profiler,nvprof
Debug工具： Nsight, CUDA-GDB, CUDA-MEMCHECK