CUDA 执行模型#

本章介绍 CUDA SIMT 执行模型——工作如何被组织为 线程、warp、block 和 grid——以及 cuda-oxide 如何通过安全、符合人体工学的 Rust API 暴露每一层级。

参见

CUDA 编程指南 -- 编程模型 是关于 CUDA 执行模型的权威参考。

线程、block 和 grid#

每次 kernel 启动都会创建一个由 thread block 组成的 grid。 这三层层次结构是 GPU 编程的基础：

层级

是什么

大小

关键属性

Grid

单次 kernel 调用启动的所有 block

每个维度最多 2³¹ - 1 个 block

Block 之间独立执行

Block

可以协作的一组线程

最多 1024 个线程

线程共享快速片上内存

Warp

block 内 32 个连续线程

固定 32

以锁步方式执行指令（SIMT）

一次 kernel 启动指定两件事：grid 中有多少个 block（grid 维度）， 以及每个 block 中有多少个线程（block 维度）。然后硬件自动将 每 32 个连续线程分组为 warp——你不需要显式创建 warp。

gpu-programming/images/simt-thread-hierarchy.svg

CUDA 三层线程层次结构。一个 2×2 的 block 组成的 grid，每个 block 包含 256 个线程，排列为 8 个各含 32 线程的 warp。底部图例将 CUDA 概念映射到 其 cuda-oxide API 等价物。#

cuda-oxide 中的线程索引#

在 kernel 内部，每个线程需要知道它应该处理哪个元素。 CUDA 提供了内置变量（threadIdx、blockIdx、blockDim、 gridDim）；cuda-oxide 将这些封装在 cuda_device::thread 模块中：

use cuda_device::{kernel, thread, DisjointSlice};

#[kernel]
pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
    let idx = thread::index_1d();
    if let Some(c_elem) = c.get_mut(idx) {
        *c_elem = a[idx.get()] + b[idx.get()];
    }
}

thread::index_1d() 计算 blockIdx.x * blockDim.x + threadIdx.x——将 每个线程映射到恰好一个数组元素的全局平面索引。这是 1D 数据并行 kernel 的 常见情况。

对于需要单独访问各分量值的情况，cuda-oxide 暴露了原始访问器：

cuda-oxide API

等价 CUDA C++

返回值

thread::index_1d()

blockIdx.x * blockDim.x + threadIdx.x

全局 1D 线程索引

thread::threadIdx_x()

threadIdx.x

线程在其 block 内的位置

thread::blockIdx_x()

blockIdx.x

Block 在 grid 中的位置

thread::blockDim_x()

blockDim.x

每个 block 的线程数（x 维度）

小技巧

对于多维索引（例如 2D 矩阵运算），结合使用 threadIdx_y()、 blockIdx_y() 和 blockDim_y() 以及 _x 变体来计算行/列索引。

Warp 和 SIMT 执行#

Warp 是 NVIDIA GPU 上的基本调度单元。一个 block 中每 32 个 连续线程组成一个 warp，且一个 warp 中的所有 32 个线程在同一时间 执行相同的指令——但作用于不同的数据。这种模型称为 SIMT（单指令、多线程）。

当一个 warp 中的所有线程遵循相同的控制流路径时，该 warp 可以达到 满吞吐量。当线程发生分支发散（不同线程走不同的 if 分支）， 硬件会串行化执行路径：它先执行一个分支，在此期间某些线程被屏蔽， 然后执行另一个分支，最后重新收敛。这被称为分支发散，它会 直接降低吞吐量。

gpu-programming/images/simt-warp-execution.svg

左图：统一执行，所有 32 个线程在一个周期内执行相同的指令。 右图：分支发散，奇数和偶数线程走不同的路径，需要两次串行执行。#

为什么这很重要#

你不需要考虑 warp 就能编写正确的 kernel——cuda-oxide 会处理这些细节。但理解 SIMT 有助于编写高效的 kernel：

  • 优先使用统一控制流。 当一个 warp 中所有线程判断相同分支条件时， 没有发散惩罚。

  • 数据相关的分支是可以的，只要相邻线程（同一 warp 内的线程）倾向于 走相同的路径。

  • 避免基于线程 ID 的分支，例如在热循环中使用 if thread::threadIdx_x() % 2 == 0——这保证每个 warp 都会发散。

参见

CUDA 编程指南 -- SIMT 架构 关于 warp 执行和重新收敛的完整硬件规范。

硬件映射#

当你启动一个 kernel 时，GPU 的硬件调度器将每个 block 分配给一个 Streaming Multiprocessor（SM）。多个 block 可以在同一个 SM 上并发运行—— 具体数量取决于 block 的资源使用情况（寄存器、 共享内存、线程数）。

关键洞察：你控制 grid 和 block 维度；硬件控制其他一切。 你永远不需要指定哪个 SM 运行哪个 block，或者 block 以什么顺序执行。 这种分离使得同一个 kernel 可以从只有少数 SM 的笔记本 GPU 扩展到 拥有 100 多个 SM 的数据中心 GPU。

gpu-programming/images/simt-hardware-mapping.svg

GPU 调度器将 8 个 block 分配给 4 个 SM。每个 SM 有自己的 warp 调度器、CUDA 核心以及共享内存/L1 缓存。Block 4-7（虚线箭头） 在 block 0-3 完成后运行，或者在资源允许时与它们一起排队运行。#

限制并发的因素#

每个 SM 有固定的资源池。只有当 SM 对以下所有资源都有足够余量时， 一个 block 才会被分配到该 SM：

资源

典型限制（Ampere）

由谁控制

线程数

每个 SM 2048 个

block_dim

寄存器

每个 SM 65536 个

编译器分配

共享内存

每个 SM 164 KB（可配置）

shared_mem_bytes

Block 槽位

每个 SM 32 个

Grid 大小

当一个 block 完成时，其资源被释放，调度器立即将一个排队的 block 分配 到该 SM。这就是为什么启动比 GPU 拥有的 SM 更多的 block 不仅是没问题的—— 而且是正常的、预期的模式。

参见

CUDA 编程指南 -- 硬件实现 关于各架构的 SM 资源限制和 occupancy 计算。

启动配置#

在主机端，LaunchConfig 告诉运行时如何构建 grid：

use cuda_core::LaunchConfig;

// 快速 1D 启动：每个 block 256 个线程，足够的 block 覆盖 N 个元素
let cfg = LaunchConfig::for_num_elems(N as u32);

for_num_elems 使用 256 的 block 大小，并通过向上取整除法计算 grid 大小——对于 大多数逐元素 kernel 来说是正确的默认值。如需更多控制， 直接构造 LaunchConfig：

let cfg = LaunchConfig {
    grid_dim: (4, 4, 1),      // 4×4 = 16 个 block
    block_dim: (16, 16, 1),   // 16×16 = 每个 block 256 个线程
    shared_mem_bytes: 0,       // 无动态共享内存
};

然后把它传给生成的启动方法：

module
    .vecadd(&stream, LaunchConfig::for_num_elems(N as u32), &a_dev, &b_dev, &mut c_dev)
    .expect("Kernel launch failed");

或者使用异步 API：

module
    .vecadd_async(LaunchConfig::for_num_elems(N as u32), &a_dev, &b_dev, &mut c_dev)?
    .sync()?;

选择 block 大小#

Block 大小是最重要的调优参数：

  • 256 线程是一个安全的默认值。它在大多数架构上平衡了 occupancy（每个 SM 多个 block）和寄存器压力。

  • 2 的幂次（128、256、512）自然地与 warp 边界对齐， 避免浪费线程。

  • 太小（< 128）可能导致 warp 调度器未充分利用。

  • 太大（1024）使用完整的 block 线程限制，可能减少 每个 SM 上并发 block 的数量。

Grid 大小由 block 大小和问题规模决定： grid_x = (N + block_x - 1) / block_x。这正是 LaunchConfig::for_num_elems 的计算逻辑。

边界检查#

因为 grid 大小是向上取整的，一些线程的索引会超出数组长度。 cuda-oxide 的 DisjointSlice 安全地处理了这种情况——get_mut 对越界索引返回 None，因此这些线程什么也不做：

#[kernel]
pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
    let idx = thread::index_1d();
    if let Some(c_elem) = c.get_mut(idx) {   // 越界线程跳过
        *c_elem = a[idx.get()] + b[idx.get()];
    }
}

这是与 CUDA C++ 有意做出的区别，在 CUDA C++ 中边界检查是程序员的责任。 cuda-oxide 的方法以一次分支（这对 warp 中除最后一个 block 外的所有线程是统一的） 为代价，消除了一整类越界内存 bug。