# CUDA 编程：从基础到应用

# 一、什么是 CUDA

CUDA 是 NVIDIA 推出的一种通用并行计算平台和编程模型，可以利用 GPU 的强大计算能力来加速各种应用程序。
CUDA 的优势在于：
- 提供了一套简单易用的编程接口，支持 C/C++/Fortran/Python 等多种语言。
- 兼容各种操作系统，如 Windows/Linux/MacOS 等。
- 支持多种 GPU 架构，如 Tesla/Fermi/Kepler/Maxwell/Pascal/Volta/Turing/Ampere 等。
- 支持多种并行编程模式，如数据并行/任务并行/流并行等。
- 支持多种优化技术，如共享内存/纹理内存/常量内存/原子操作/同步机制等。
Why CUDA?
- 串行速度提升已经结束
  - 无法继续提升频率
  - 难以继续降低功耗
- 当前计算机性能提升趋势
  - 计算机没有变得更快，而是变得更宽
    - 多核 CPU、GPU、超级计算机
  - 数据级别并行
    - 同样的指令作用于多个数据
  - 线程级别的并行

# 二、CPU vs. GPU

CPU 和 GPU 都是计算机中的重要组件，但它们有着不同的设计目标和特点。
CPU 的特点是：
- 拥有较少的核心数，但每个核心都有较高的时钟频率和较强的运算能力。
- 拥有较大的缓存和复杂的控制流机制，可以有效地降低延迟和提高串行代码的性能。
- 更适合于处理复杂的单任务或少量的多任务，如操作系统/数据库/编译器等。
- 类比于摩托车，可以灵活地在城市中穿梭。
GPU 的特点是：
- 拥有较多的核心数，但每个核心都有较低的时钟频率和较弱的运算能力。
- 拥有较小的缓存和简单的控制流机制，可以有效地提高吞吐量和利用大规模并行架构。
- 更适合于处理大量相似或简单的任务，如图形渲染/科学计算/机器学习等。
- 类比于大巴车，可以承载更多的乘客。

	CPU	GPU
缓存	大缓存：掩盖较长的存储器延迟	小缓存：但通过更快的存储提高吞吐量
运算器	强大的运算器：降低运算延迟	更节能的运算器：延迟大但总吞吐量大
控制机制	复杂的控制机制：分支预测等	简单的控制流机制：无分支预测
线程	线程高度轻量级：大量并发	线程高度轻量级：大量并发

# 三、异构计算

异构计算是指利用不同类型的处理器协同工作来完成一个任务，如 CPU+GPU、CPU+FPGA、CPU+ASIC 等。
异构计算的优势在于：
- 可以充分发挥每种处理器的特长，提高性能和效率。
- 可以降低功耗和成本，延长设备寿命和节约资源。
- 可以增加灵活性和可扩展性，适应不同场景和需求。
异构计算的挑战在于：
- 需要设计合适的编程模型和接口，实现不同处理器之间的协调和通信。
- 需要考虑不同处理器之间的负载均衡和数据一致性，避免性能瓶颈和错误发生。
- 需要优化不同处理器之间的数据传输和转换，减少开销和延迟。
CPU+GPU
- 利用 CPU 处理复杂控制流
- 利用 GPU 处理大规模运算
- CPU 与 GPU 之间通过 PCIe 总线通信
  - 新显卡支持 NVLink 连接（5-12 倍 PCIe3.0)

# 四、CUDA 编程模型

CUDA 编程模型是基于数据并行思想设计的一种分层抽象模型，可以将一个复杂的问题分解为多个简单的子问题，并将其映射到 GPU 上执行。
CUDA 编程模型包括以下几个层次：
- 线程（thread）：线程是 CUDA 中最基本的执行单元，每个线程都有自己独立的寄存器、指令指针、栈空间等。
- 块（block）：块是由多个线程组成的一维或二维的逻辑分组，每个块都有自己独立的共享内存、同步机制等。
- 网格（grid）：网格是由多个块组成的一维或二维的逻辑分组，每个网格都有自己独立的全局内存、常量内存、纹理内存等。
- 设备（device）：设备是指 GPU 本身，包括多个流式处理器（SM）、多个 CUDA 核心（core）、多个缓存、总线等。
- 主机（host）：主机是指 CPU 本身，包括内存、硬盘、键盘、鼠标等。
CUDA 编程模型的执行流程如下：
- 在主机端编写并行代码，称为核函数（kernel），并使用__global__修饰符标记。
- 在主机端调用核函数，并使用<<< grid,block >>>语法指定网格和块的维度，称为执行配置。
- 在设备端执行核函数，每个块被分配到一个 SM 上，每个线程被分配到一个 core 上。
- 在设备端完成核函数后，返回主机端继续执行后续代码。
2 级架构
- 每个 GPU 拥有多个 Streaming Multiprocessor(SM)
  - 具体数目及设计因产品而异
  - SM 共用显存
- 每个 SM 拥有多个 CUDA core
  - 数目因产品而异
  - Core 共用调度器和指令缓存
2 级架构下的执行模型：线程束（warp）
- CUDA 线程以 32 个为一组在 GPU 上执行
  - 线程束以单指令多线程的方式运行（SIMT）
    - 所有线程在不同数据上执行相同的指令
- SMIT、SIMD、SMT
  - 灵活度：SIMD < SIMT < SMT
  - 性能： SIMD > SIMT > SMT
  - SIMT 与 SIMD 相比：多个状态寄存器，多个地址，独立的执行路径
- SM 负责调度并执行线程束
  - 线程束调度时会产生上下文切换
  - 调度方式因架构而异
Host 与 device
- Host（CPU）相关：运行在 CPU 上的代码及主机内存
- Device（GPU）相关：运行在 GPU 上的代码及显存（设备内存）
- 通过在主机上调用核函数（kernel）执行并行代码
指明 host 与 device 代码
- __host__从主机端调用，在主机端执行
- __global__从主机端调用，在设备端执行
- __device__从设备端调用，在设备端执行
- __host__和__device__可以一起使用
- <<< 1,4 >>>：执行配置
  - 指明网格中有 1 个块
  - 每块中有 4 个线程
- cudaDeviceSynchronize()
  - 与 OpenMP 不同，CUDA 核函数为异步执行
- 核函数限制条件（__global__函数）
  - 只能访问设备内存
  - 必须返回 void
  - 不支持可变数量的参数
  - 参数不可为引用类型
  - 不支持静态变量
指明网格及块的维度
- 形式为<<< grid,block >>>
  - grid 与 block 为 dim3 类型
  - grid 与 block 的大小受到计算能力的限制
- GPU 架构与线程执行
  - 一个 CUDA core 执行一个线程
  - 一个 SM 执行一个 block 中的线程
  - GPU 中执行 grid 中的所有线程
- 确定线程编号
  - 使用内置变量 threadIdx、blockIdx、blockDim
- CUDA 编程例子：向量加法
```
__global__ void VecAdd(int *a, int *b, int *c)
{
  int tid = blockIdx.x;
  if (tid < N) {
    c[tid] = a[tid] + b[tid];
  }
}
int main()
{
    ...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
    ...
}
```
GPU 内存管理：
- 创建：cudaMalloc
- 拷贝：cudaMemcpy
  - 使用 cudaMemcpyHostToDevice 与 cudaMemcpyDeviceToHost 指明拷贝方向
- 释放：cudaFree
处理错误
- 使用宏定义
block 中最大线程限制：n 必须不大于 1024
同一个 block 只在一个 SM 上执行：没有充分利用 GPU 计算资源
思路：使用多个 block
- 每个 block 使用 m 个 thread（如 m=32）
- grid,block 设置：<<< n/m, m >>>
  - n 无法被 m 整除？
  - 需对 n/m 向上取整
  - 需判断 tid 是否会超过范围
- 确定 thread 的全局编号

# 五、CUDA 线程执行模型

逻辑视图
- 每个线程块由一个 SM 执行
- 由硬件调度
- 无法控制线程块的执行顺序
硬件视图
- 所有线程块在硬件上都是一维的
- 三维线程将沿 x->y->z 的顺序展开到一维
- 展开后的一维线程每 32 个形成一个线程束
  - 最后不足 32 的部分也将创建线程
    - 不活跃
    - 仍将消耗 SM 资源
线程束调度
- 线程束切换开销为 0
  - SM 保存每个线程束的执行上下文
  - 在整个线程束的生命周期中保存于芯片内
  - 上下文切换没有损失
  - 可切换同一 SM 上不同线程块的线程束
- SM 中常驻线程块数量受可用资源限制
  - 资源：程序计数器、寄存器、共享内存
  - 活跃线程束：具备计算资源的线程束
    - Kepler 上最大为 64
    - 选定的线程束：被调度到执行单元的线程束（Kepler 上最大为 4）
    - 符合条件的线程束：准备执行但尚未执行
    - 阻塞的线程束：没做好执行准备（指令参数未就绪，无可用 CUDA core）
- 活跃线程束于延迟隐藏
  - 满载：线程调度器在每个时钟周期都有符合条件的线程束
  - 通过调度符合条件的线程束，可以有效的掩盖指令延迟
    - 算数指令：算数操作从开始到产生输出（10~20 时钟周期）
    - 内存指令：发出加载/存储操作到数据到达目的地（全局内存~800 时钟周期）
  - 应适当增加活跃线程束
    - Little’s law
    - 线程数不宜过少（每个线程处理的任务数与线程数需要平衡）
    - 线程块资源不易过多（如，共享内存的大小与活跃线程块数量需要平衡）
线程束执行
- 每个线程束以 SIMD 方式在 SM 上执行
  - 线程束内同时执行同样语句
  - 线程束外的视角看来为 SIMT
- 分支分化
  - 线程束出现不同的控制流
  - 性能优化：避免分支分化，因为线程束只能执行相同的逻辑，在执行某一个路径的线程时会禁用另一路径的线程。
- busy waiting vs signal
  - busy waiting：如，使用 while 循环不断检查条件是否满足
  - signal：当条件满足由系统发送指令
    - __syncthreads()：只能在线程块内同步，不能在不同的线程块同步
  - busy waiting 的问题：死锁
- 减少分支分化的影响
  - 减少 if 语句
    - 尤其是减少基于 threadIdx 的 if 语句
    - 使用条件赋值代替条件语句
  - 平衡分支执行时间
    - 避免出现执行时间过长的分支

# 六、CUDA 原子操作

原子指令
- 执行过程不能分解为更小的部分：不被中断
- 避免竞争条件出现
- 竞争条件
  - 程序运行结果依赖于不可控的执行顺序
CUDA 原子操作：
- 基本操作：atomicCAS
  - 其它所有原子操作均可由 atomicCAS()实现
  - CAS：compare and swap
    - 读取目标位置(address)并与预期值(old_val)进行比较
      - 相等则将 new_val 写入目标位置
      - 不相等则不发生变化
    - 返回目标位置中原值：可用来检查 CAS 操作是否成功
原子指令与并发控制：原子指令在并发控制中起着重要的作用。在多线程或多进程的环境中，当多个线程或进程尝试同时访问和修改共享数据时，如果没有适当的控制机制，可能会导致数据的不一致性。原子指令通过确保某些操作在执行过程中不会被其他线程或进程中断，来避免这种情况。
竞争条件与死锁：竞争条件是并发编程中的一个主要问题，它发生在两个或更多的线程或进程在无序或未同步的情况下访问和修改共享数据，导致结果不可预测。原子指令是解决竞争条件的一种方法，但也可能引入另一个问题 - 死锁。死锁是指两个或更多的进程或线程互相等待对方释放资源，导致所有进程或线程都无法继续执行。
CUDA 原子操作与 GPU 编程：在 GPU 编程中，由于大量的线程并行执行，可能会有多个线程同时访问和修改同一块内存。CUDA 的原子操作提供了一种机制，使得在这种情况下仍能保证数据的一致性。然而，过度依赖原子操作可能会导致性能下降，因为它们违反了 GPU 编程的基本原则——并行执行。因此，在设计 GPU 算法时，应尽量减少原子操作的使用，或者寻找可以避免使用原子操作的算法。
CUDA 原子操作的应用：在某些情况下，CUDA 原子操作是必要的。例如，在统计或计数问题中，需要多个线程共享一个计数器，并且每个线程都可能需要增加计数器的值。在这种情况下，使用 CUDA 原子操作可以保证计数器的正确性。另一个例子是图形处理，其中可能需要多个线程同时更新像素的值。使用 CUDA 原子操作可以避免同时更新导致的数据不一致问题。