CUDA编程中的GPU执行模型概述

CUDA编程中的GPU执行模型概述

引用

CSDN

1.

https://m.blog.csdn.net/qq_42681630/article/details/144908366

CUDA编程中的GPU执行模型是理解并行计算的关键。本文将深入探讨GPU架构、SM（流式多处理器）、线程级别并行、指令级并行、线程束及其执行特点等内容，帮助读者全面了解CUDA编程中的GPU执行机制。

GPU架构概述

GPU由多个流式多处理器（Streaming Multiprocessors, SM）组成，每个SM又包含多个CUDA核心（CUDA Cores）和其他硬件资源。GPU中每个SM都能支持数百个线程并发执行，每个GPU通常有多个SM，当一个核函数的网格被启动的时候，多个block会被同时分配给可用的SM上执行。

在SM上：

• 同一个线程块（Thread Block）内，多个线程同时执行，实现并行计算。
• 在单个线程内，通过流水线技术同时执行多条指令，提高指令执行效率。

以下是一个SM的架构图：

补充

1. 线程级别并行（Thread-Level Parallelism, TLP）

• 定义：线程级别并行是指在同一时间内，多个线程同时执行不同的任务。
• 实现方式：
• 在GPU中，一个线程块（Thread Block）内的所有线程会被分配到同一个SM上执行。
• SM通过硬件资源（如CUDA核心、共享内存等）同时管理多个线程的执行。
• 线程块内的线程可以协作完成任务，例如通过共享内存交换数据。
• 特点：
• 线程之间是独立的，每个线程可以执行不同的代码路径。
• 线程块内的线程数量通常较大（例如256或512个线程），以实现高并行度。
• 示例：
• 在矩阵加法中，每个线程负责计算一个矩阵元素，所有线程同时执行，实现线程级别并行。

2. 指令级并行（Instruction-Level Parallelism, ILP）

• 定义：指令级并行是指在单个线程内，通过流水线技术同时执行多条指令，以提高指令执行效率。
• 实现方式：
• 现代处理器（包括GPU的CUDA核心）使用流水线技术将指令执行分为多个阶段（如取指、解码、执行、写回等）。
• 在同一时间内，流水线的不同阶段可以处理不同的指令，从而实现指令级并行。
• 例如，当一个指令在执行阶段时，下一个指令可以在解码阶段，而第三个指令可以在取指阶段。
• 特点：
• 指令级并行是硬件自动实现的，程序员无需显式管理。
• 通过提高指令吞吐量，减少单个线程的执行时间。
• 示例：
• 在计算a = b + c; d = e * f;时，加法指令和乘法指令可以同时进入流水线的不同阶段，实现指令级并行。

线程束

在CUDA编程中，线程束（Warp）是GPU执行并行计算的基本单位。

• 线程束是GPU中一组32个线程的集合，是SM（流式多处理器）调度和执行的最小单位。一个SM可以同时执行多个线程束。
• SIMT架构：CUDA采用单指令多线程（SIMT, Single Instruction Multiple Threads）架构
• 单指令（Single Instruction）：所有线程执行相同的指令。
• 多线程（Multiple Threads）：多个线程同时执行，每个线程可以处理不同的数据。

线程束的执行

• 同步执行：线程束中的所有线程在同一个时钟周期内执行相同的指令。这种同步执行使得线程束能够高效地利用GPU的计算资源。
• 发散（Divergence）：当线程束中的线程遇到条件分支（如if语句）时，可能会导致不同线程执行不同的路径。这种情况称为线程束发散。发散会导致性能下降，因为不同路径的指令需要顺序执行，直到所有路径都完成。
• 合并（Convergence）：当所有分支路径执行完毕后，线程束会重新合并，继续同步执行

SIMT的发散（Divergence）

• 如果线程束中的线程执行不同的分支（例如if-else语句），GPU会串行执行这些分支，导致性能下降
• 例如：

if (threadIdx.x % 2 == 0) {
    // 分支 1
} else {
    // 分支 2
}

• 在这个例子中，线程束中的线程会分成两组（偶数和奇数），分别执行不同的分支。
• GPU会先执行分支1，再执行分支2，导致性能下降。

SIMT和SIMD的区别

两者都是将相同指令广播给多个执行单元，但是SIMT的某些线程可以选择不执行，也就是说同一时刻所有线程被分配给相同的指令，SIMD规定所有人必须执行，而SIMT则规定有些人可以根据需要不执行，这样SIMT就保证了线程级别的并行，而SIMD更像是指令级别的并行。

Fermi 架构

最早的架构

Fermi架构是NVIDIA的经典GPU架构之一，具有以下核心特性：

• 加速核心：512个CUDA核，每个CUDA核包含一个全流水线的整数算术逻辑单元(ALU)和一个浮点数运算单元(FPU)。
• SM组织：CUDA核被组织到16个SM（流式多处理器）上。
• 内存接口：6个384-bit的GDDR5内存接口，支持6GB的全局机载内存。
• 线程管理：GigaThread引擎负责分配线程块到SM的线程束调度器上。
• 二级缓存：768KB的二级缓存，被所有SM共享。

每个SM包括以下资源：

• 执行单元：CUDA核，负责执行整数和浮点运算。
• 调度单元：线程束调度器和指令调度单元，负责调度线程束的执行。
• 存储资源：共享内存、寄存器文件和一级缓存，用于快速数据访问和线程通信。

每个SM的具体执行细节如下：

• 加载/存储单元：每个SM有16个加载/存储单元，因此每个时钟周期内有16个线程（半个线程束）可以计算源地址和目的地址。
• 特殊功能单元(SFU)：执行固有指令，如正弦、余弦、平方根和插值。SFU在每个时钟周期内的每个线程上执行一个固有指令。
• 线程束调度：每个SM有两个线程束调度器和两个指令调度单元。当一个线程块被分配到SM时，线程块内的所有线程被分成线程束。两个线程束调度器选择两个线程束，并通过指令调度器存储它们要执行的指令。
• 线程束切换：在同一个SM上，线程束之间的切换由硬件自动管理，且没有时间开销。当一个线程束等待内存访问或其他操作时，SM会立即切换到另一个就绪的线程束，以充分利用计算资源。

Fermi架构还支持并发执行核函数，允许多个核函数同时在GPU上执行，从而提高GPU的利用率和任务并行性。

Kepler 架构

Kepler是Fermi的提升版本，首先是硬件上的加强

优化1：内核可以启动内核

可以通过GPU完成递归操作

优化2：hyper-Q

Hyper-Q技术主要是CPU和GPU之间的同步硬件连接，以确保CPU在GPU执行的同事做更多的工作。Fermi架构下CPU控制GPU只有一个队列，Kepler架构下可以通过Hyper-Q技术实现多个队列如下图。