【读书笔记】【CUDA编程指南】CUDA简介

第1章 CUDA介绍——深入理解CUDA编程模型

本文基于 CUDA Programming Guide Release 13.2 第1章整理，覆盖 GPU 硬件架构、线程执行模型、内存体系以及 CUDA 编译链的完整知识体系，适合有一定 C++ 基础、希望入门 GPU 编程的读者。

目录

• 第1章 CUDA介绍——深入理解CUDA编程模型
  
  
  
  • 目录
    
  • 1. GPU 计算的崛起
    
  • 2. GPU vs CPU：设计哲学的根本差异
    
  • 3. 快速开始：使用 GPU 的多种途径
    
    
    
    • 使用现成库
      
    • 使用 AI 框架
      
    • 使用领域专用语言（DSL）
      
    
    
  • 4. CUDA 编程模型
    
    
    
    • 4.1 异构系统：Host 与 Device
      
    • 4.2 GPU 硬件：SM 与 GPC
      
    • 4.3 线程层次结构
      
      
      
      • 内建变量
        
      • 线程块的执行规则
        
      • 线程束（Warp）
        
      • Cluster（簇）—— 计算能力 9.0 新增
        
      
      
    • 4.4 SIMT 执行模型与分支发散
      
      
      
      • 分支发散（Warp Divergence）
        
      • SIMT vs SIMD
        
      
      
    
    
  • 5. GPU 内存体系
    
    
    
    • 5.1 全局内存（DRAM）
      
    • 5.2 片上内存
      
    • 5.3 统一内存
      
    
    
  • 6. CUDA 平台与编译链
    
    
    
    • 6.1 计算能力（Compute Capability）
      
    • 6.2 驱动、Toolkit、Runtime 与 Driver API 的关系
      
    • 6.3 PTX 虚拟 ISA
      
    • 6.4 cubin、fatbin 与编译产物
      
    • 6.5 二进制兼容性与 PTX 兼容性
      
      
      
      • cubin 的二进制兼容性规则
        
      • PTX 的前向兼容性
        
      
      
    • 6.6 JIT 编译
      
    
    
  • 7. 总结
    
  
  

1. GPU 计算的崛起

GPU（图形处理器）最初是为实时 3D 渲染设计的专用定制硬件，负责固定功能的并行图形流水线加速。到了 2003 年，图形渲染管线的部分阶段变得完全可编程，开发者可以为场景中的每一个像素或顶点运行自定义代码。
这一可编程能力引发了研究者的想象：能否把 GPU 强大的并行计算能力用到图形以外的领域？2006 年，NVIDIA 正式推出 CUDA（Compute Unified Device Architecture，统一计算设备架构），使得开发者可以脱离图形 API，将任意计算工作负载交给 GPU 来加速。
自此，GPU 计算被广泛用于：

• 流体动力学、能量传输等科学仿真
  
• 数据库与大数据分析
  
• 图像分类、扩散模型、大语言模型等 AI 技术
  

GPU 的可编程性已成为现代 AI 基础设施不可或缺的基石。
2. GPU vs CPU：设计哲学的根本差异

GPU 和 CPU 在设计目标上有着根本差异。
维度CPUGPU核心目标最快执行单条串行线程（低延迟）同时执行数千线程（高吞吐）并发线程数几十个数万至数十万个晶体管分配大量用于缓存和流控制大量用于数据计算单元单线程性能强相对弱总计算吞吐量相对低极高下图直观展示了两者的晶体管分配差异：

CPU 以缓存和流控制（Control）为主，GPU 则将绝大多数晶体管投入到计算核心（ALU）。

这一差异决定了 GPU 非常适合数据并行（Data-Parallel）的计算场景：大量数据做同样的操作，而不需要复杂的串行依赖关系。
3. 快速开始：使用 GPU 的多种途径

直接编写 CUDA 代码并不是使用 GPU 的唯一方式。根据场景不同，有以下几种由浅入深的途径：
使用现成库

NVIDIA 提供了大量针对各 GPU 架构深度优化的库，往往比自己实现的性能更好：
库名用途cuBLAS线性代数（矩阵乘法等）cuFFT快速傅里叶变换cuDNN深度神经网络CUTLASS模板化矩阵乘法使用 AI 框架

PyTorch、TensorFlow 等框架在底层调用上述库来实现 GPU 加速，用户无需感知 CUDA 细节。
使用领域专用语言（DSL）

• NVIDIA Warp：面向物理仿真的 Python DSL，编译至 CUDA 执行。
  
• OpenAI Triton：面向深度学习算子开发的 Python DSL，编译至 PTX/CUDA。
  

以上途径均构建于 CUDA 平台之上。本文后续聚焦于直接编写 CUDA 代码的核心概念。
4. CUDA 编程模型

4.1 异构系统：Host 与 Device

CUDA 编程模型假设一个异构计算系统，即同时包含 CPU 和 GPU 的系统。
术语含义HostCPU 及其直接连接的内存（主机内存）DeviceGPU 及其直接连接的显存（设备内存）Host Code运行在 CPU 上的代码Device Code / Kernel运行在 GPU 上的代码CUDA 应用始终从 CPU 开始执行。Host Code 通过 CUDA API 完成三件事：

• 在 Host 内存和 Device 内存之间拷贝数据
  
• 启动（Launch）核函数（Kernel）在 GPU 上并行执行
  
• 等待数据拷贝或 GPU 执行完成
  

CPU 和 GPU 可以同时运行，最优性能通常来自最大化两者的利用率。
4.2 GPU 硬件：SM 与 GPC

对于 CUDA 编程而言，GPU 可以被理解为以下硬件层次：

1. GPU
2. └── GPC（图形处理集群，Graphics Processing Cluster）×N
3.     └── SM（流多处理器，Streaming Multiprocessor）×M
4.         ├── 本地寄存器文件（Register File）
5.         ├── 统一数据缓存（Unified Data Cache）
6.         │   ├── L1 Cache（自动缓存）
7.         │   └── Shared Memory（可编程共享内存）
8.         └── 多个功能计算单元（ALU 等）

复制代码

每个 SM 内部的 Unified Data Cache 在物理上是同一块资源，但可以在运行时配置 L1 Cache 和 Shared Memory 的比例分配。
下图展示了 GPU 的完整硬件架构：

GPU 通过 PCIe 或 NVLINK 与 CPU 相连。

4.3 线程层次结构

CUDA 的线程组织是理解编程模型的核心。从小到大依次为：

1. Thread（线程）
2.     ↓ ×32
3. Warp（线程束）
4.     ↓ ×N
5. Thread Block（线程块）
6.     ↓ ×M [可选，计算能力 9.0+]
7. Cluster（簇）
8.     ↓ ×K
9. Grid（网格）

复制代码

每启动一个 Kernel，整体执行的所有线程构成一个 Grid（网格）。

内建变量

在 Kernel 内，每个线程通过以下内建变量确定自己的"身份"：
变量类型含义threadIdxdim3线程在其线程块内的位置（x/y/z）blockIdxdim3线程块在网格内的位置（x/y/z）blockDimdim3线程块的维度大小（x/y/z）gridDimdim3网格的维度大小（x/y/z）线程块和网格均支持 1、2、3 维布局，方便将线程映射到不同形状的数据。
以一维网格为例，计算每个线程的全局唯一索引：

1. __global__ void kernel() {
2.     int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;
3.     // 使用 globalIdx 决定该线程负责处理哪个数据元素
4. }

复制代码

线程块的执行规则

• 同一线程块内的所有线程保证运行在同一个 SM 上，可通过共享内存高效通信与同步。
  
• 不同线程块之间相互独立，执行顺序任意，不能有数据依赖关系。
  
• 一个 Grid 可以包含数百万个线程块，GPU 的调度器将其分发至可用的 SM 执行。
  

下图展示了 Grid 中的线程块如何被分配到各 SM 上执行：

每个 SM 同时承载多个线程块，但线程块之间的调度顺序没有任何保证。

线程束（Warp）

线程块内的线程以 32 个为一组组成一个 Warp（线程束）。Warp 是 GPU 调度和执行的基本单元。

最佳实践：线程块的线程总数最好是 32 的整数倍。若不是整数倍，最后一个 Warp 会有若干空闲 Lane，导致计算单元浪费。

Cluster（簇）—— 计算能力 9.0 新增

随着 GPU 内 SM 数量的不断增多，跨线程块的通信如果只能依赖全局内存，效率极低。为此，计算能力 9.0（Hopper 架构，如 H100）引入了簇（Cluster）这一可选分组层级，位于线程块和网格之间。
簇的特性：

• 同一簇内的所有线程块调度到同一 GPC 内的各 SM 上，且同时执行。
  
• 簇内不同块的线程可通过 Cooperative Groups API 进行通信与同步。
  
• 簇内线程可访问同 GPC 内所有线程块的共享内存，称为分布式共享内存（Distributed Shared Memory）。
  

下图展示了含簇的网格逻辑结构（Grid → Cluster → Block 三层嵌套）：

下图展示了簇内线程块如何被同时调度到同一 GPC 内的各 SM：

4.4 SIMT 执行模型与分支发散

Warp 内 32 个线程以 SIMT（Single Instruction Multiple Threads，单指令多线程） 方式执行：所有线程同时执行同一条指令，但每个线程操作的数据不同（由各自的 threadIdx/blockIdx 决定）。
分支发散（Warp Divergence）

当 Warp 内的线程走上不同的分支时，就会发生分支发散：

1. if (threadIdx.x % 2 == 0) {
2.     a = r(t);   // 仅偶数 lane 执行
3. } else {
4.     a = q(t);   // 仅奇数 lane 执行
5. }
6. y = f(a);       // 全部 lane 执行

复制代码

执行过程：

• 先执行 if 分支，奇数 lane 被掩码屏蔽（masked off）等待
  
• 再执行 else 分支，偶数 lane 被掩码屏蔽等待
  
• 最后所有 lane 执行 y = f(a)
  

这意味着两段代码串行执行，原本并行的能力被折半浪费。下图直观展示了各 Lane 在分支执行时的掩码状态：

绿色 Lane 为活跃执行，灰色 Lane 为掩码等待状态。因此，减少 Warp 内分支发散是 CUDA 性能优化的重要目标。

SIMT vs SIMD

SIMDSIMT控制流统一，单一路径每个线程可走独立路径数据宽度固定（如 AVX-512 = 512 bit）无固定宽度编程抽象向量寄存器独立线程SIMT 给程序员提供了"每个线程独立"的编程抽象，但底层仍是 Warp 粒度的执行。

注意：指南明确建议程序员按照 SIMT 编程模型编写代码，而非依据底层硬件实现细节。硬件在不违反编程模型的前提下可进行优化；若程序依赖未定义的 Warp 执行细节，可能在不同 GPU 架构上产生 UB（未定义行为）。

5. GPU 内存体系

现代计算中，内存的高效利用与计算单元的最大化同等重要。GPU 拥有多种类型的内存，各有其用途和访问特性。
5.1 全局内存（DRAM）

• GPU 自带的 DRAM 称为全局内存（Global Memory），即通常所说的"显存"，所有 SM 均可访问。
  
• CPU 的 DRAM 称为系统内存/主机内存（System Memory / Host Memory）。
  
• 在当前支持的系统中，CPU 和 GPU 共享同一虚拟地址空间，但各自的物理内存相互独立（每个 GPU 的虚拟地址范围唯一且与 CPU 不同）。
  
• CUDA 提供 API（如 cudaMalloc、cudaMemcpy）用于：在 GPU 上分配内存、在 CPU 和 GPU 之间或 GPU 之间进行数据拷贝。
  

5.2 片上内存

除全局内存外，每个 SM 还有多种片上内存（On-Chip Memory），访问速度极快：
内存类型归属访问范围控制方式说明寄存器文件（Register File）每个 SM单个线程私有编译器自动分配存储线程局部变量，速度最快共享内存（Shared Memory）每个 SM同一线程块（或簇）内全部线程程序员显式控制线程间通信的高速通道L1 Cache每个 SM自动（透明）硬件自动管理与共享内存共享物理资源（Unified Data Cache）L2 Cache整个 GPU所有 SM 共享硬件自动管理比 L1 更大，访问稍慢常量缓存（Constant Cache）每个 SMKernel 内所有线程只读声明为 __constant__缓存全局内存常量及 Kernel 参数关键资源约束：调度线程块到 SM 时，该线程块所需的寄存器总数（每线程寄存器数 × 线程数）不能超过 SM 的寄存器文件大小；超出则 Kernel 无法启动，需减少线程块大小或每线程寄存器用量。
5.3 统一内存

统一内存（Unified Memory） 是 CUDA 提供的一项特性，允许用 cudaMallocManaged 分配一块 CPU 和 GPU 均可访问的内存。CUDA 运行时或底层硬件负责在需要时将数据迁移到正确的设备。

1. float *data;
2. cudaMallocManaged(&data, N * sizeof(float));
3. // CPU 和 GPU Kernel 均可直接访问 data，无需手动 cudaMemcpy

复制代码

即使使用统一内存，最优性能的原则仍然是：尽量让数据待在最常访问它的设备上，减少内存迁移次数。

另外还有一种映射内存（Mapped Memory），是将 CPU 内存映射为 GPU 可直接访问的地址。但由于访问要跨越 PCIe/NVLINK 总线，延迟高、带宽低，GPU 无法用并行度掩盖这一开销，性能远不如统一内存，一般不推荐作为数据访问的主要方式。
6. CUDA 平台与编译链

6.1 计算能力（Compute Capability）

每块 NVIDIA GPU 都有一个计算能力（Compute Capability，CC）编号，格式为 X.Y（主要版本号.次要版本号），表示该 GPU 所支持的特性集以及若干硬件参数。

• 计算能力与 SM 版本直接对应：CC 12.0 的 GPU，其 SM 版本为 sm_120。
  
• 可在 CUDA GPU Compute Capability 页面 查询所有 GPU 的计算能力。
  

6.2 驱动、Toolkit、Runtime 与 Driver API 的关系

这是初学者最容易混淆的部分，以下分层说明：

1. ┌──────────────────────────────────────────────┐
2. │          你的 CUDA 应用                       │
3. ├──────────────────────────────────────────────┤
4. │          CUDA 运行时 API（cudaXxx）            │  ← CUDA Toolkit 的一部分
5. ├──────────────────────────────────────────────┤
6. │          CUDA 驱动 API（cuXxx）               │  ← 由 NVIDIA 驱动暴露
7. ├──────────────────────────────────────────────┤
8. │          NVIDIA 驱动（r580 等版本）            │  ← GPU 的"操作系统"
9. ├──────────────────────────────────────────────┤
10. │          GPU 硬件                             │
11. └──────────────────────────────────────────────┘

复制代码

组件说明NVIDIA 驱动GPU 的操作系统，所有 GPU 用途（包括显示、Vulkan、DirectX）的基础，版本号如 r580CUDA Toolkit独立软件包，包含头文件、库（cuBLAS 等）和工具（nvcc、nvprof 等）CUDA 运行时（CUDA Runtime）Toolkit 中的核心库，提供 cudaMalloc、cudaMemcpy、Kernel Launch 等高层 APICUDA 驱动 API由 NVIDIA 驱动暴露的底层 API（函数以 cu 开头），功能更灵活，部分高级特性只有它才有运行时 API vs 驱动 API：运行时 API 是驱动 API 的高层封装，使用更方便；驱动 API 更底层、更灵活。两者可以在同一应用中混用。
6.3 PTX 虚拟 ISA

PTX（Parallel Thread Execution，并行线程执行） 是 CUDA 平台中往往"隐身"的重要一层——它是 NVIDIA GPU 的虚拟指令集架构（Virtual ISA），可以理解为一种面向 GPU 的高级汇编语言。
PTX 为真实 GPU 的物理指令集提供了一层抽象，使得：

• 编译器（如 nvcc）先将 C++ 代码编译为 PTX（中间表达），再将 PTX 编译为特定 GPU 的二进制。
  
• 其他语言或 DSL（如 Triton）可以生成 PTX，借助 NVIDIA 的工具链完成最终编译。
  
• 运行时可通过 NVRTC 库将 CUDA C++ 在线动态编译为 PTX，适合需要生成动态 GPU 代码的场景。
  

PTX 版本与计算能力对应，例如支持 CC 12.0 所有特性的 PTX 版本称为 compute_120。
6.4 cubin、fatbin 与编译产物

术语含义cubinPTX 编译后的 GPU 本地二进制，针对特定 SM 版本（如 sm_120）fatbin一个容器，内部可同时打包多个版本的 cubin 和/或 PTX一个可执行文件或库同时包含 CPU 二进制代码和 fatbin 容器，运行时根据当前 GPU 从 fatbin 中选取最适合的 cubin 加载执行。

6.5 二进制兼容性与 PTX 兼容性

cubin 的二进制兼容性规则

在同一主要版本的计算能力下：

• GPU 可以加载次要版本号小于等于自身的 cubin。
  

例如，GPU 计算能力为 12.5：
cubin 版本能否加载原因sm_115❌主要版本不同sm_120✅主要版本相同，次要版本 0 ≤ 5sm_125✅主要版本相同，次要版本 5 ≤ 5sm_129❌次要版本 9 > 5sm_90❌主要版本不同

二进制兼容性仅对 nvcc 等 NVIDIA 官方工具生成的二进制有效；手动修改二进制将使兼容性承诺失效。

PTX 的前向兼容性

PTX 可在运行时被驱动 JIT 编译为等于或高于其目标计算能力的任意 cubin。例如，fatbin 中包含 compute_80 的 PTX，可在运行时编译为 sm_120 的 cubin。
这使得应用无需重新编译即可在将来更新的 GPU 架构上运行。
6.6 JIT 编译

JIT（Just-in-Time，即时编译） 指在程序运行时由 NVIDIA 驱动将 PTX 编译为本地 cubin 的过程。

• 优点：前向兼容未来 GPU；每次升级驱动可自动享受编译器优化改进。
  
• 缺点：首次运行会增加程序启动时间。
  
• 缓存机制：JIT 编译的结果会被缓存（称为 compute cache），驱动升级后缓存自动失效以触发重新编译。
  

7. 总结

本章涵盖了理解 CUDA 编程的基础框架：
主题要点GPU 设计哲学高吞吐量，数千线程并行，牺牲单线程性能异构编程Host (CPU) 负责调度，Device (GPU) 负责并行计算线程层次Thread → Warp(32) → Block → Cluster(CC9.0+) → Grid内建变量threadIdx/blockIdx/blockDim/gridDimSIMT同 Warp 内同指令异数据，分支发散会降低效率内存层次寄存器 > 共享内存/L1 > L2 > 全局内存，速度依次递减计算能力决定可用特性，格式 X.Y，与 SM 版本一一对应编译链C++ → PTX → cubin，打包进 fatbin，JIT 提供前向兼容理解上述概念是高效编写 CUDA 程序的前提。后续章节将围绕这些概念展开具体的 C++ API 使用和性能优化实践。
参考资料：NVIDIA CUDA Programming Guide Release 13.2, Chapter 1

来源：程序园用户自行投稿发布，如果侵权，请联系站长删除
免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！