从CUDA编程的角度看GPU结构

显卡结构

显卡的结构设计始终围绕 “并行效率” 和 “数据吞吐”：

1. 整体层面：通过 “GPU + 显存 + 供电 + 散热” 的协同，确保硬件能稳定输出高算力；
2. GPU 内部：通过 “GPC→SM→计算单元” 的分层架构，将海量任务拆解为并行子任务，最大化利用 CUDA/Tensor/RT 核心；
3. 软硬衔接：通过 “Block→Warp” 的硬件调度单位，将软件定义的并行线程映射到实际计算单元，避免资源闲置。

网格结构

网格中的块会被 GPU 的全局调度器（如 GigaThread Engine）分配到多个 SM 上执行（一个 SM 可同时运行多个块，只要资源足够）。例如，100 个块可能被分配到 10 个 SM 上，每个 SM 处理 10 个块。网格大小可远超 GPU 的 SM 数量（如 10000 个块），未被调度的块会在队列中等待，适合处理大规模任务（如百万级数据的并行计算）。

• SM 不绑定任何网格，而是通过调度器动态处理块。
• 多个网格的块可被分配到同一 SM，但每个块始终属于一个特定的网格。
• 网格是核函数调用的产物，SM 的硬件设计不感知网格的存在，仅关注块的资源占用和执行效率。

SM（流式多处理器）

SM（流式多处理器）的最小执行单元是线程束（Warp），而非单个线程。

• 线程束的定义：一个线程束是固定数量的线程集合（NVIDIA GPU 中通常为 32 个线程，AMD 为 64 个，这里以 NVIDIA 为例）。线程束内的所有线程执行相同的指令流（单指令多数据，SIMD），但操作不同的数据（数据并行）。
• 调度逻辑：SM 的 Warp 调度器（Warp Scheduler）只能以线程束为单位分配计算资源（如 CUDA 核心、Tensor Core），无法单独调度单个线程。即使核函数中线程的逻辑分支不同（如if-else），线程束内的线程也会通过 “掩码执行”（Masked Execution）保持指令同步（不满足条件的线程暂时闲置）。

SM的硬件设计会感知线程束（Warp）和线程块（Block）的存在，但不感知网格（Grid）和单个线程（Thread）

GPU 的设计采用 “软件定义并行粒度，硬件优化执行效率” 的分层模式：

• 软件层（网格、块、线程）：程序员通过网格划分任务规模、通过块组织线程协作、通过线程定义最小计算单元，实现对并行任务的灵活拆分（如 1000 个块 ×256 个线程 = 256000 线程的并行规模）。
• 硬件层（线程束）：SM 通过固定大小的线程束实现高效并行执行，硬件优化线程束的调度（隐藏延迟）、资源利用（如寄存器 Bank 冲突避免）和指令同步（掩码执行），将软件定义的灵活并行转换为硬件可高效处理的固定模式。

这种分层设计既保留了软件的灵活性（可适配从几千到几亿线程的任务），又确保了硬件的执行效率（通过固定大小的线程束简化调度和电路设计）。

1. SM 对线程束（Warp）的强感知：硬件直接管理的执行单元

线程束是 SM 硬件设计的 “原生单位”，SM 通过专用硬件组件直接控制其执行：

• Warp 调度器（硬件组件）：每个 SM 包含多个 Warp 调度器（如 Ada 架构为 4 个），负责从就绪的线程束中选择指令、分发到计算核心（CUDA 核心 / Tensor 核心），并管理线程束的状态（如 “运行中”“等待内存”）。
• 指令同步机制：线程束内的 32 个线程必须执行相同指令（SIMT 模型），硬件通过 “掩码寄存器” 标记活跃线程（如if-else分支中不满足条件的线程被掩码屏蔽），这种同步逻辑由硬件电路直接实现。
• 资源绑定：线程束的执行依赖 SM 的物理资源（如寄存器端口、共享内存 Bank），硬件会为每个线程束分配固定的资源切片，确保并行执行时无冲突。

简言之，线程束是 SM 硬件的 “基本语言”，所有计算任务最终都必须转换为线程束的执行流程，SM 的硬件设计完全围绕线程束的调度和执行展开。

2. SM 对线程块（Block）的感知：资源分配与执行边界

线程块是软件定义的逻辑单位，但 SM 必须感知其存在以实现资源管理：

• 资源分配单元：每个块在被调度到 SM 前，SM 会根据块的需求（如请求的寄存器数量、共享内存大小）检查自身剩余资源，只有资源充足时才会接收该块（这一过程由 SM 的资源分配器硬件完成）。
• 块内同步支持：SM 提供__syncthreads()的硬件实现 —— 当块内线程调用该函数时，SM 会暂停整个块的执行，直到所有线程到达同步点（通过块级别的计数器硬件实现）。
• 执行隔离：一个块的所有线程束只能在分配的 SM 上执行（块不能跨 SM 拆分），SM 通过硬件标记区分不同块的资源（如共享内存分区、寄存器空间），确保块之间的资源隔离。

因此，SM 虽不参与块的定义，但必须感知块的存在以实现资源分配、同步和隔离，块是 SM 与软件逻辑之间的 “接口单位”。

3. SM 对网格（Grid）和线程（Thread）的 “无感知”

• 网格（Grid）：网格是多个块的集合，由全局调度器（如 NVIDIA 的 GigaThread Engine）管理，SM 仅接收全局调度器分配的单个块，完全不关心该块属于哪个网格，也不参与网格级的资源协调。
• 线程（Thread）：单个线程是软件逻辑的最小单位，但 SM 的硬件设计中没有 “线程级” 的调度或资源分配 —— 线程的执行完全依赖其所属的线程束，硬件通过线程束内的索引（如threadIdx.x）区分线程，而非单独管理每个线程。