从 CUDA 代码到 GPU 执行：Grid、Block、CTA、Warp 与 SM 如何对应

发表于 2026-10-24 分类于 GPU 阅读次数：本文字数： 12k 阅读时长 ≈ 44 分钟

一个 kernel 启动时，程序员给出的只有两个数字：多少个 Block，每个 Block 多少线程。这两个数字怎样变成 SM 上真实发生的执行，取决于几层固定的分组规则——Block 内的线程被切成 Warp，Warp 由 SM 调度，Grid 里的 Block 分批进入硬件。

先用一次向量加法建立编程视角，再看硬件怎样执行，最后回到同一个例子把对应关系走完。Cluster、CLC、Stream 会沿这条线依次出现，分别改变协作范围、工作分配和启动顺序。

一、编程视角：线程怎样组织

1. Kernel 与 Thread

要算 c[i] = a[i] + b[i]，两个数组各长 1,000。CPU 上是一个循环，CUDA 上可以让每个线程处理一个元素。

Kernel 是在 GPU 上执行的函数，Thread 是这份代码的一个执行实例。每个线程靠自己的编号决定处理哪份数据，并持有自己的寄存器状态和执行位置。线程数可以远多于计算单元数，硬件会安排它们分批推进。

核心代码如下，显存分配、数据传输和错误检查都省略，d_a、d_b、d_c 指向已分配的设备内存：

#include <cuda_runtime.h>

__global__ void add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

void launch_add(const float* d_a, const float* d_b, float* d_c, int n) {
    if (n <= 0) return;
    constexpr int threads_per_block = 256;
    int blocks = (n + threads_per_block - 1) / threads_per_block;
    add<<<blocks, threads_per_block>>>(d_a, d_b, d_c, n);
}

__global__ 声明 kernel，<<<blocks, threads_per_block>>> 是执行配置：启动多少个线程块、每块多少线程。它没有说明占用几个 SM。启动对 CPU 通常异步，读取计算结果前需要同步。CUDA C++ 入门

`<<< >>>` 的完整四参数形式

普通 CUDA C++ kernel 启动可以写成下面的形式。尖括号内是执行配置，圆括号内才是传给 kernel 函数的参数：

1	kernel<<<grid_dim, block_dim, shared_mem_bytes, stream>>>(...);

配置项	含义	能否省略
第 1 项：Grid 维度	`dim3`，指定各方向的 Block 数量	必须提供
第 2 项：Block 维度	`dim3`，指定各方向的线程数量	必须提供
第 3 项：动态共享内存	`size_t`，每个 Block 额外申请的共享内存字节数	可省略，默认 `0`
第 4 项：Stream	`cudaStream_t`，本次启动提交到哪个 CUDA 流	可省略，默认 `0`，即默认流

第三项的单位是字节／Block，与 kernel 内静态声明的共享内存另计。例如第三项为 1,024，表示每块额外申请 1 KB；kernel 内可通过 extern __shared__ 数组使用这块内存。它不表示每线程 1 KB，也不是所有 Block 合计 1 KB。CUDA 执行配置定义

用前面的 add 给出一个四项全部写出的例子，设备内存与错误检查仍按前例省略：

void launch_add_full(const float* d_a, const float* d_b, float* d_c, int n) {
    if (n <= 0) return;
    dim3 block_dim(256, 1, 1);  // 每块 256 × 1 × 1 个线程
    dim3 grid_dim((n + 255) / 256, 1, 1);

    size_t shared_mem_bytes = 0;  // 这个 add 不需要动态共享内存
    cudaStream_t stream;
    cudaStreamCreate(&stream);

    add<<<grid_dim, block_dim, shared_mem_bytes, stream>>>(d_a, d_b, d_c, n);

    cudaStreamSynchronize(stream);  // 本示例等待该流的工作完成
    cudaStreamDestroy(stream);
}

Stream 可以先理解为 GPU 工作的提交队列：同一个流中的工作按提交顺序执行，不同流之间的工作在条件允许时可以重叠，不能仅凭创建了多个流就认定一定并行。实际程序通常会复用已有的流，不必每次启动都创建、同步和销毁；完整的顺序与并发规则放在第三章。CUDA Streams

常见的省略方式如下，smem 表示需要的动态共享内存字节数：

add<<<grid_dim, block_dim>>>(d_a, d_b, d_c, n);           // 第 3、4 项均为 0
add<<<grid_dim, block_dim, 0>>>(d_a, d_b, d_c, n);        // 省略第 4 项
add<<<grid_dim, block_dim, 0, stream>>>(d_a, d_b, d_c, n); // 无动态共享内存，指定流
// 若某个 kernel 需要动态共享内存，其配置可写成 <<<grid_dim, block_dim, smem, stream>>>

要指定第 4 项，必须保留第 3 项的位置。<<<grid_dim, block_dim, stream>>> 不表示“省略共享内存、直接指定流”，第三个位置仍会被当成共享内存字节数，通常直接编译失败。

前两项可以使用 dim3 表达一维、二维或三维。写 256 相当于 dim3(256,1,1)；写 dim3(16,16) 时，省略的 z 默认是 1，每块仍有 256 个线程。例如 Grid 为 dim3(4,3) 时共有 12 个 Block。使用二维、三维配置时，kernel 也需要相应处理 y、z 索引，不能直接套用上面只处理 x 方向的 add。

Cluster 等更高级的运行时启动属性可通过 cudaLaunchKernelEx 配置，不是在普通 <<< >>> 中继续追加第 5 项。

2. Block 与 CTA

Block（Thread Block，线程块）是一组能通过共享内存和块内屏障协作的线程；CTA（Cooperative Thread Array，协作线程数组）是同一件事在 PTX 里的叫法。CUDA C++ 代码写 Block，PTX、CUTLASS 和架构资料写 CTA，两者之间没有第三层。

同一个 Block 内，线程用 __shared__ 声明共享内存，用 __syncthreads() 等块内线程到达同一屏障，并让屏障之前的内存操作对块内可见。典型用法是先合作搬一块数据，同步，再一起使用。屏障要求块内所有线程参与，不能只放进一部分线程会走的分支里。CUDA 块内同步

Block 可以是一维、二维或三维。dim3(16,16) 的 256 个线程便于对应图像坐标，256 是一维写法。维度服务于编号方式，硬件并不会照着这个形状摆放线程。

3. Grid

一次 kernel 启动产生的线程块集合称为 Grid（网格）。同一 Grid 里所有 Block 的线程数和形状相同。

代码里的四个内建变量分别回答不同的问题：

变量	含义	本文示例：`n = 1000`
`threadIdx.x`	线程在块内的横向编号	0～255
`blockDim.x`	每个块横向有多少线程	256
`blockIdx.x`	块在 Grid 内的横向编号	0～3
`gridDim.x`	Grid 横向有多少块	4

于是 Block 0 负责元素 0～255，Block 1 负责 256～511，Block 2 负责 512～767，Block 3 负责 768～999。总共 1,024 个线程，多出的 24 个线程被 i < n 挡掉，不读写数组。Block 2 里编号为 17 的线程，全局元素编号是 2 × 256 + 17 = 529。这条公式把“线程在块内的位置”和“块在 Grid 中的位置”拼成了工作编号。

Grid 可以开得很大，同时在跑的 Block 数量由硬件资源决定。普通 kernel 不保证 Block 的执行顺序：Block 3 可能先于 Block 0，某些 Block 也可能还在排队等资源。程序不能假设编号靠前的块先完成，也不能把块内屏障当成整个 Grid 的屏障。跨块同步要靠专门机制，或者拆成多次 kernel 启动。CUDA 编程模型

4. Cluster：把协作范围扩大到多个 Block

Hopper 引入了 Thread Block Cluster（线程块集群），在 Grid 和 Block 之间插入一层可选分组。同一个 Cluster 的 Block 保证在同一个 GPC（Graphics Processing Cluster，图形处理集群）内协同调度，可以互相访问共享内存，这种跨块访问称为分布式共享内存（Distributed Shared Memory）。Hopper 调优指南



%%{init: {"flowchart": {"nodeSpacing": 22, "rankSpacing": 28, "padding": 12}, "themeVariables": {"fontSize": "15px"}}}%%
flowchart TB
    grid["Grid<br/>一次启动的全部线程块"]
    cluster["Cluster · 可选<br/>多个 Block 协作"]
    block["Block / CTA<br/>组内共享内存与同步"]
    thread["Thread<br/>执行 kernel 的线程实例"]
    grid -->|启用 Cluster| cluster
    cluster -->|包含| block
    grid -->|普通启动| block
    block -->|包含| thread

图 1：编程层级。两条分支对应两种组织方式，本文的向量加法走普通启动。

Cluster 不抹掉 Block 的边界：每个 Block 仍有自己的线程编号、共享内存和资源需求，多出来的是跨 Block 的协作能力。这里要区分三个名字：Cluster 是线程块分组，GPC 是硬件组织，CLC 是调度控制机制。GPC 与 Cluster 分别属于硬件侧和编程侧，一个 Cluster 不会跨 GPC。

二、硬件视角：线程在哪里执行

5. SM

SM（Streaming Multiprocessor，流式多处理器）是 GPU 中执行线程的主要硬件单元。一块 GPU 有很多 SM，若干 SM 组成一个 GPC——Cluster 的协同调度范围就是 GPC，这也是上一节那条约束的来源。中间的其他层级和图形专用单元这里不展开。

一个 SM 大致由这些资源组成：

SM 内的资源	主要作用
Warp Scheduler	从就绪 Warp 中选择指令发射
CUDA Core 等算术单元	执行相应的浮点、整数等运算
Tensor Core	加速支持的数据类型和形状下的矩阵乘加
Load/Store 单元	执行访存相关指令
寄存器文件	保存线程的寄存器状态
共享内存／L1	提供片上数据存储和缓存

CUDA Core 是运算资源，不是一个能独立调度线程的 CPU 核心；线程也不永久绑定在某个 CUDA Core 上，指令按类型进入不同管线。上面的向量加法不会因为跑在带 Tensor Core 的 GPU 上就自动变成 Tensor Core 计算。CUDA 硬件模型、NVIDIA 对 Tensor Core 的说明

6. CTA 如何落到 SM 上

一个 CTA 的线程由同一个 SM 执行，一个 SM 可以同时驻留多个 CTA。普通 CTA 不会拆成两半，分别放到两个 SM 上跑。

驻留指 CTA 的线程状态和所需资源已经分配到 SM，并不表示它的所有线程在同一个时钟周期执行同一条指令。SM 内部还要继续挑选哪些线程组可以推进。



%%{init: {"flowchart": {"nodeSpacing": 20, "rankSpacing": 28, "padding": 12, "subGraphTitleMargin": {"top": 12, "bottom": 20}}, "themeVariables": {"fontSize": "15px"}}}%%
flowchart TB
    pending["待调度 CTA<br/>资源足够时进入 SM"]
    subgraph sm["一个 SM · 简化示意"]
        a["驻留 CTA A<br/>包含多个 Warp"]
        b["驻留 CTA B<br/>包含多个 Warp"]
        scheduler["Warp Scheduler<br/>选择就绪 Warp 的指令"]
        units["计算／访存管线<br/>执行发射的指令"]
        a --> scheduler
        b --> scheduler
        scheduler --> units
    end
    pending -.-> a

图 2：CTA 分配和指令发射是两个粒度。图中的两块只是某一时刻的可能状态，既不表示 SM 固定驻留两块，也不表示调度器数量。

能驻留几个 CTA，受寄存器、共享内存、线程数、Warp 数和 Block 数上限共同限制。每块占用大量共享内存时，即使算术单元还空着，也可能放不下第二块；每线程的寄存器需求升高，可驻留的线程数同样下降。CUDA 硬件多线程

所以 <<<4, 256>>> 不能翻译成“四个 SM，每个 SM 跑 256 个线程”。四个 CTA 落在哪里，由硬件调度和当时的资源决定。

7. Warp：SM 怎样组织线程

NVIDIA GPU 把同一个 CTA 内的线程按每 32 个一组切成 Warp（线程束），Lane 是线程在 Warp 内的位置，编号 0～31。Warp 只由同一 CTA 的线程组成，不会把两个 CTA 的零散线程拼在一起。

一维的 256 线程 Block 划分如下：

Warp 0线程 0～31

Warp 1线程 32～63

Warp 2线程 64～95

Warp 3线程 96～127

Warp 4线程 128～159

Warp 5线程 160～191

Warp 6线程 192～223

Warp 7线程 224～255

图 3：一个 CTA 包含 8 个 Warp，每个 Warp 都有自己的 Lane 0～31，编号属于当前 Block。

二维、三维 Block 先按 x 最快、y 次之、z 最慢展平成线性编号，再每 32 个一组。Warp 数等于线程数除以 32 后向上取整：33 个线程需要两个 Warp，第二个只有一个有效线程。CUDA 线程层级

一条指令怎样覆盖 32 个线程

从程序看，32 个线程各自执行同一份 kernel、各自处理数据；从硬件看，它们被组织成一个 Warp，尽量一起执行同一条指令。程序决定 Block 多大，硬件按上面的编号规则切出 Warp。

32 个线程都执行 x = x * 2，每个线程的 x 可以不同。硬件发射一条乘法指令，参与的线程各自相乘，不必为每个线程单独发一份相同指令。下一条加法或访存指令同样如此，调度器还可以在中间插入其他就绪的 Warp。

这就是 SIMT（Single Instruction, Multiple Threads，单指令多线程）。共同执行指令不要求数据相同，也不要求结果相同：向量加法里，一个 Warp 的线程读不同的 a[i]、b[i]，做同一个加法，写不同的 c[i]。相邻线程访问相邻元素，还便于合并全局内存请求。CUDA 的 Warp 与 SIMT 模型

分支分歧：一个 Warp 里的两条路径

分支条件由每个线程单独判断，一个线程成立不代表同一 Warp 的其他线程也成立。

下面这个 kernel 只启动一个 Block、32 个线程，恰好是一个 Warp：Lane 0～15 走 A 段，Lane 16～31 走 B 段。

// 仅用于 branch_demo<<<1, 32>>>(d_in, d_out)
__global__ void branch_demo(const float* in, float* out) {
    int lane = threadIdx.x;
    float x = in[lane];

    if (lane < 16) {
        x = x * 2;  // A 段
        x = x + 1;
    } else {
        x = x * 3;  // B 段
        x = x - 1;
    }

    out[lane] = x;
}

先假定 A、B 被编译成两段不同的指令，编译优化放到后面再说。发射 A 段的某条指令时，硬件要确定哪些线程参与；B 段是另一条指令，不能并进同一次发射。

这个参与状态叫活跃线程掩码（Active Mask）。可以把它想成 32 个开关：某个 Lane 的开关打开就参与当前指令，关闭就不执行，也不会因此误写自己的结果。执行 A 段时只有前 16 个打开，执行 B 段时只有后 16 个。

执行阶段	Lane 0～15 16 个线程	Lane 16～31 16 个线程
读取 x	参与 · 读取自己的值	参与 · 读取自己的值
执行 A 段	参与 · 乘 2，再加 1	不参与 A 段
执行 B 段	不参与 B 段	参与 · 乘 3，再减 1
写回结果	参与 · 写入自己的结果	参与 · 写入自己的结果

图 4：同一 Warp 内两条分支的参与情况。“不参与”只针对当前指令，不表示线程退出。这里只画一种便于理解的次序，A、B 的实际先后和交错方式没有保证。

同一个 Warp 的线程走了不同路径，称为 Warp Divergence（线程束分歧）。分歧不改变成员：原来的 32 个线程仍在原来的 Warp 里，也不会和别的 Warp 中走相同分支的线程重新组队。到了后面共同的代码位置，线程可以重新一起执行，这个过程叫重新汇合（Reconvergence）。SIMT 分支执行说明

分歧的代价

一条 Warp 指令理想情况下有 32 个线程参与。上例执行 A 或 B 时都只有 16 个，而两条路径都得处理。结果是对的，但一部分指令只服务了一半线程。

设 A、B 各有 100 条需要分别发射的指令，忽略公共代码和控制开销：这个 Warp 要处理约 200 条路径指令，每条 16 个线程参与，两段的平均 Lane 参与率是 50%。这是个便于理解的指令计数模型，不表示整个 kernel 一定慢一倍——实际耗时还受访存、指令类型、编译优化和其他 Warp 的调度影响。

循环会有同类问题。同一 Warp 里大多线程只转一圈、少数要转很多圈时，后面的循环指令就只有少量线程参与。

但这不意味着 CUDA 里不能写 if。是否分歧，取决于同一 Warp 内的判断结果是否一致。一个一维 Block 有 64 个线程，前 32 个组成 Warp 0，后 32 个组成 Warp 1：

条件	Warp 0 的选择	Warp 1 的选择	是否出现 Warp 内分歧
`threadIdx.x < 32`	全部走 A	全部走 B	两个 Warp 内都没有
`threadIdx.x % 2 == 0`	各一半走 A、B	各一半走 A、B	两个 Warp 内都有

不同 Warp 各走各的路径没有问题。优化时该看的是某个 Warp 是否长期只有少数线程参与，而不是数源码里有多少个 if。CUDA 控制流最佳实践

另外，源码里的 if/else 不一定编译成跳转。分支很短时，编译器可能改用谓词执行（Predication）：给指令附一个执行条件，条件不成立的线程跳过这条指令。上面几行算术就可能被这样处理，也可能被合并。图示解释的是参与关系，不能直接推出最终的机器指令和运行时间。

同属一个 Warp，仍然需要同步

分歧决定每个线程做什么，同步解决一个线程什么时候能安全使用另一个线程的结果，两个问题要分开处理。

从 Volta 开始，独立线程调度让硬件按线程维护执行位置等状态，推进和重新汇合可以更细。SIMT 的执行方式没变，分歧仍可能浪费发射机会；程序也不能假定 Warp 内 32 个线程永远逐条同步前进。

每个线程先往共享内存写一个值、再读相邻线程的值，就必须保证对方已经写完：

// 仅用于 warp_neighbor<<<1, 32>>>(d_in, d_out)
__global__ void warp_neighbor(const float* in, float* out) {
    __shared__ float values[32];
    int lane = threadIdx.x;
    values[lane] = in[lane];

    __syncwarp(0xffffffffu);  // 32 个线程都到达；使此前写入对彼此可见

    out[lane] = values[lane ^ 1];  // 0 读 1，1 读 0，2 读 3，3 读 2……
}

0xffffffffu 是 32 位全 1 的参与掩码，表示这次同步覆盖整个 Warp。它是程序声明的参与范围，不是某一刻硬件的活跃线程状态。这个例子里 32 个线程都会执行到同步点；只有部分线程参与时，要按 API 规则给出正确的掩码，不能把全掩码搬进分支。

寄存器之间的线程交换还可以用 __shfl_sync 这类 Warp 集合操作；跨 Warp 共享数据，通常需要块内同步这类范围更大的机制。属于同一个 Warp 只是分组事实，数据依赖仍然要靠程序自己同步。NVIDIA 的 Warp 同步与集合操作说明

Warp Specialization：让不同 Warp 承担不同职责

Warp 0 整体走 A、Warp 1 整体走 B 时，两个 Warp 内部都没有分歧。Warp Specialization（线程束专门化，也叫 Warp 分工）把这种格局变成长期安排：不同 Warp 固定承担不同职责。它是 kernel 的一种写法，没有新增线程层级，也没有把 Warp 拆小。

常见分工是一组 Warp 从显存搬数据，另一组用已经到达的数据计算。前者叫 Producer（生产者），后者叫 Consumer（消费者），生产的是可供计算使用的数据。实际实现也可以按多个 Warp 组成的组来分。CUDA 的生产者／消费者 Warp 分工示例

一个 128 线程的 Block 有 4 个 Warp，可以安排 Warp 0 搬运、Warp 1～3 计算。下面只演示角色选择，省掉了实际工作和同步，不能直接运行：

int warp_id = threadIdx.x / 32;

if (warp_id == 0) {
    // Producer：等待空缓冲区，发起数据搬运，通知数据已准备好
} else {
    // Consumer：等待数据，执行计算，通知缓冲区可以重用
}

Warp 0 里 32 个线程算出的 warp_id 都是 0，整体进入生产者路径；其余三个 Warp 也各自整体进入消费者路径。角色差异落在 Warp 之间，选择本身没有制造 Warp 内分歧。若改成按线程奇偶分工，同一个 Warp 里就会同时有生产者和消费者，执行特征和同步设计都要另算。

某个 Warp 负责搬运，也不代表它的 32 个线程一直在搬。TMA 可以由少量线程发起大块传输，剩下的交给异步硬件。分工描述的是职责归属，换算不出 Lane 的利用率。Hopper 的 TMA 与 Warp 分工

分工换来的是重叠

每次都等搬运完成再开始计算，两段耗时要相加。用异步搬运加两个共享内存缓冲区，可以让生产者准备下一份数据的同时，消费者算上一份：

阶段	Producer Warp 0	Consumers Warp 1～3
启动	加载 tile 0 写入缓冲区 0	等待首份数据
流水阶段 1	加载 tile 1 写入缓冲区 1	计算 tile 0 读取缓冲区 0
流水阶段 2	加载 tile 2 重用缓冲区 0	计算 tile 1 读取缓冲区 1
流水阶段 3	加载 tile 3 重用缓冲区 1	计算 tile 2 读取缓冲区 0
收尾	加载已完成	计算 tile 3 读取缓冲区 1

图 5：一个双缓冲安排。阶段表示依赖次序，不表示固定时长；同一行的搬运与计算在资源允许时可以重叠。缓冲区重用前，必须等上一份数据使用完毕。

重叠发生在“加载 tile 1”和“计算 tile 0”之间。计算 tile 0 仍要等 tile 0 到达，数据依赖没有消失；双缓冲让读写落在不同区域，避免生产者覆盖消费者正在用的数据。

流水线进入稳态后，处理下一份工作的间隔接近搬运和计算中较慢的那一项，不再是两者相加。但启动要先把数据备好，收尾还要算完剩下的，小任务不一定受益。分工也不是唯一做法：同一组 Warp 照样可以设计异步多级流水线，分工只是组织重叠的一种方式。CUTLASS 的流水线与 Warp 分工设计

分工还能让各组 Warp 的工作更集中：搬运组管地址和传输，计算组管矩阵运算，后处理组管结果变换和写回。它们仍然共享同一个 SM 的资源，分工并不会额外分配硬件。

分工与分歧的区别

对比项	Warp Divergence	Warp Specialization
描述什么	同一 Warp 内的线程走不同路径	不同 Warp／Warp 组承担不同职责
典型例子	一个 Warp 内按线程奇偶执行 A、B	Warp 0 搬运，Warp 1～3 计算
对执行的影响	不同路径分别推进，部分指令只有部分线程参与	不同 Warp 按各自角色推进，可组织搬运与计算重叠
性能关注点	路径长度、每条指令的参与线程数	各阶段吞吐、等待时间、缓冲区和资源分配
可能的代价	持续分歧使 Lane 参与率降低	分工不均或同步开销抵消重叠收益

两者可以同时出现在一个 kernel 里。按 Warp 分好生产者、消费者之后，消费者内部仍可能因为处理边界元素出现分歧。分工不会自动消除分歧；分歧也不说明写错了，要看实际指令、工作量和等待情况。

Warp 分工常见的性能问题有：搬运跟不上，计算组等数据；计算跟不上，搬运组等空缓冲区；搬运组分到太多 Warp，挤占计算资源；共享内存缓冲区开得太多，能同时驻留的 CTA 变少。数据量很小或本来就有充分重叠时，新增分工只剩下管理成本。

正确性落在两次交接上：写完才能读，用完才能覆盖。发起异步搬运不等于搬运已经完成；消费者发出异步矩阵计算后，也要确认相关数据不再被使用，才能释放缓冲区。

这些交接一般用 pipeline、异步屏障表达。__syncwarp() 只管指定 Warp 内部，不能让生产者 Warp 等到别的消费者 Warp。若在只有生产者进入的分支里放一个要求整个 Block 参与的 __syncthreads()，而尚未退出的消费者不会到达该同步点，程序就会挂住。真实流水线要准确设置参与者、完成条件和缓冲区状态。CUTLASS 对 Warp 分工与同步的说明

延迟隐藏：换一个就绪 Warp 发射

某个 Warp 等待访存结果时，下一条相关指令暂时发不出去。Warp Scheduler 可以挑另一个就绪 Warp 发射，让计算管线继续工作。GPU 能这样做，是因为它保留了大量线程状态。

这里涉及三个状态：驻留指 Warp 的执行状态已经在 SM 上；就绪指当前可以发射某条指令；发射指指令已被选中并送进相应管线。驻留多，就绪未必多——很多 Warp 可能都在等数据。

32 个 Lane 也不能理解成“占住 32 个 CUDA Core，在一个周期里做完所有指令”。吞吐和延迟取决于管线和具体架构。

8. Occupancy：线程更多不等于更快

Block 大小影响 Warp 数和资源分配：太大，一个 SM 可能只能驻留少数几块；太小，容易先撞上驻留 Block 数上限。线程总数也不能无限提高有限管线和显存带宽的吞吐。

Occupancy（占用率）通常指 SM 上驻留 Warp 数占该设备允许最大驻留 Warp 数的比例。它提供的是候选 Warp 的多少，不直接等于算术单元利用率：所有 Warp 都在等待时，占用率可以很高而指令发射依然不足；反过来，有些 kernel 在较低占用率下就已经打满吞吐或带宽。CUDA 最佳实践：占用率

选启动配置要同时看任务划分、寄存器、共享内存、访存和计算瓶颈。Grid 与 Block 决定程序表达出多少并行，CTA 与 Warp 决定这些并行怎样分组，SM 的资源和调度决定它们推进多快。

9. CLC：运行中的 CTA 领取后续工作

前面几节已经够解释普通 kernel。CLC（Cluster Launch Control，集群启动控制）让 kernel 自己参与工作分配：运行中的 CTA 可以尝试取消一份尚未启动的工作，拿到它的索引，然后接着把那份工作做完。这一机制由 Blackwell 引入，CUDA 指南以计算能力 10.0 为起点介绍，使用时仍要核对目标设备的支持范围。CLC 官方说明

常规做法有两种。普通 kernel 每个 CTA 处理一份工作然后退出，硬件能动态分配，代价是反复做 CTA 初始化；持久化 kernel 只留少量 CTA 常驻、循环处理多份工作，初始化和状态可以复用，但后续工作怎么分配得自己安排。

普通调度本身就能在 SM 有空闲资源时补进新 CTA。CLC 多出来的是：已经跑起来的 CTA 也能接手后续工作，在复用状态的同时做动态负载均衡。

把矩阵输出切成很多 tile。某个 CTA 做完 tile 0，可以尝试领取尚未启动的 tile 7；领取成功时，原本对应 tile 7 的启动被取消，当前 CTA 用拿到的索引去算 tile 7。变的是程序使用的工作索引，当前 CTA 的 blockIdx 不会变成 7。CUTLASS 的 CLC 调度说明



%%{init: {"flowchart": {"nodeSpacing": 20, "rankSpacing": 26, "padding": 12}, "themeVariables": {"fontSize": "15px"}}}%%
flowchart TB
    first["用初始 blockIdx<br/>完成第一份工作"]
    request["请求取消尚未启动的工作<br/>等待并读取 CLC 结果"]
    success{"领取成功？"}
    next["使用取得的索引<br/>完成下一份工作"]
    finish["结束领取循环<br/>完成收尾并退出"]
    first --> request --> success
    success -->|是| next
    next --> request
    success -->|否| finish

图 6：以单 CTA 的 Cluster 为例。失败可能是没有待领取的工作，也可能是调度等其他原因；未被取消的工作仍由正常调度处理。

用流程伪代码表示：

work_id = 初始 blockIdx
完成必要的初始化
循环：
    协作处理 work_id
    发出 CLC 请求，等待结果，并向参与线程发布结果
    如果请求失败：退出循环
    work_id = 领取到的索引
完成收尾

真实实现涉及异步结果、共享内存屏障和 CTA 内的协作，这段伪代码不能当 CUDA API 用。CLC 按 Cluster 粒度工作：单 CTA 的 Cluster 对应领一个线程块，多 CTA 的 Cluster 要成组协调。它领的是尚未启动的工作，不会把另一个正在执行的 CTA 搬到当前 SM。

三、启动之间：Stream 与 Event

Grid 描述一次启动里的线程块；多次启动之间怎样排序、能不能重叠，由 Stream 管。Stream、Grid、SM 三者描述的事情不同：

概念	组织什么	与其他概念的关系
Stream	kernel 启动、内存拷贝、事件等操作的顺序	可以先后提交多个 Grid，也可以包含非计算操作
Grid	某一次 kernel 启动的全部 CTA	CTA 根据资源情况分配到 SM
SM	已分配 CTA 的执行状态和计算资源	执行其中的 Warp 指令

例如，同一个 add 函数分别启动两次，会产生两个 Grid；它们可以提交到同一个 Stream，也可以提交到不同 Stream。Stream 不属于 Grid → Block → Thread 这条线程包含关系，它管理的是一次次启动之间的关系。

普通 Stream 不绑定某个 SM，也不自动划分一份专属算力。一个 Grid 的 CTA 可以分布在多个 SM 上；来自不同 Stream、不同 Grid 的 CTA，在资源和调度允许时也可以驻留在同一个 SM。具体放在哪里、何时运行，由硬件和运行时安排。NVIDIA 对跨 kernel CTA 共驻留的说明

同一条流内：按提交顺序执行

先看同一 GPU、同一 CUDA 上下文内，通过普通 <<< >>> 提交的 kernel；本节末尾再说一种高级例外。

向 s1 依次提交 A、B、C。CPU 通常很快提交完就继续做别的事；GPU 上是 A 完成后才执行 B，B 完成后才执行 C。即使 A 暂时只用了部分 SM，B 也不能因为还有空闲资源就越过这条顺序。

这保证 B 能用上 A 写出的结果，不必在两次启动之间插一次 CPU 等待。顺序执行也不等于 GPU 一次只跑一个线程：A 自己的 CTA 和 Warp 仍然并行。只有一个 Stream 的程序，照样能同时用上多个 SM。CUDA 的流顺序语义

不同流之间：允许重叠，不保证重叠

若 s1 里有 A、B，s2 里有互不相关的 C，A、B 仍分先后，C 则可能和 A 或 B 重叠，也可能先等一会儿。CPU 先提交 B、后提交 C，并不会建立“B 完成后 C 才能跑”的跨流依赖。

Stream	较早的时间段	较晚的时间段
s1	执行 A	A 完成后执行 B
s2	独立的 C 可以与 A、B 的执行重叠

图 7：一种可能的并发安排。A、B 属于同一流，按普通启动规则不能彼此重叠；C 的开始、结束与重叠范围仅为示意，不代表保证或测量结果。

能不能重叠，取决于设备是否支持并发 kernel、还有多少 SM／寄存器／共享内存可用，以及操作之间是否存在同步依赖。看到重叠也不一定更快：两个 kernel 可能争抢显存带宽、缓存或计算管线。一个 kernel 已经把瓶颈资源占满时，多开几个 Stream 不会提高总吞吐。CUDA 并发 kernel 执行说明

用前面的 add 演示提交，假设设备数组都已分配、输入就绪、输出互不重叠，grid_dim、block_dim 与前例相同：

cudaStream_t s1, s2;
cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);

// A、B 在 s1 中按顺序执行；B 使用 A 的输出 d_c
add<<<grid_dim, block_dim, 0, s1>>>(d_a, d_b, d_c, n); // A
add<<<grid_dim, block_dim, 0, s1>>>(d_c, d_b, d_e, n); // B

// C 使用独立的数据，有机会与 A 或 B 重叠
add<<<grid_dim, block_dim, 0, s2>>>(d_u, d_v, d_w, n); // C

// 先提交所有工作，再等待；避免每次启动后立刻等待
cudaStreamSynchronize(s1);
cudaStreamSynchronize(s2);
cudaStreamDestroy(s1);
cudaStreamDestroy(s2);

cudaStreamNonBlocking 关系到默认流的同步规则，后面会讲；它不保证两个 kernel 并行，也不改变同一 Stream 内的顺序。

跨流有依赖时用 Event

如果把 B 挪到 s2，而 B 要读 A 写出的 d_c，就不能让两者无约束地重叠。CPU 先写下 A 的启动语句，不代表 A 已经执行完。可以在两个 Stream 之间用 Event（事件）挂一个依赖。

保留上例的流创建和清理，把中间的提交换成：

cudaEvent_t ready;
cudaEventCreateWithFlags(&ready, cudaEventDisableTiming);

add<<<grid_dim, block_dim, 0, s1>>>(d_a, d_b, d_c, n); // A
cudaEventRecord(ready, s1); // 标记 s1 中 A 之后的完成位置

cudaStreamWaitEvent(s2, ready, 0); // s2 后续工作等待该事件
add<<<grid_dim, block_dim, 0, s2>>>(d_c, d_b, d_e, n); // B

cudaStreamSynchronize(s2);
cudaEventDestroy(ready);

cudaStreamWaitEvent 把依赖交给 GPU 的工作队列，通常不阻塞 CPU；它约束的是 s2 中随后提交的工作，之前已经排好的独立工作仍可先跑。相比之下，cudaStreamSynchronize(s1) 让 CPU 等指定流此前提交的工作，cudaDeviceSynchronize() 等当前设备、当前上下文里此前提交的所有流。CUDA 显式同步

默认流让并发失效

省略第 4 项或填 0 就是用默认 Stream。默认流有两种模式。

Legacy default stream 与普通的 blocking stream 之间存在隐式同步。在两个 blocking stream 的任务中间插一次默认流启动，就可能把这些本来互不相干的任务串成一条线。Per-thread default stream 让每个 CPU 线程有自己的默认流，按普通流处理，可以用 nvcc --default-stream per-thread 之类的方式启用。

cudaStreamCreate() 创建的流默认参与上述 legacy 规则；cudaStreamCreateWithFlags(..., cudaStreamNonBlocking) 可以避开这类隐式同步。这里的 blocking／non-blocking 说的是默认流之间的同步关系，与“这次 kernel 启动是否阻塞 CPU”无关。默认流与 non-blocking streams

排查“为什么没有并发”时先看四件事：工作是不是都放进了同一个流，两次提交之间是不是立刻同步了，有没有混用 legacy 默认流，硬件资源够不够同时容纳两个 kernel。要确认真的重叠，得看 GPU 时间线；CPU 提交得快、或者创建了两个流，都不能说明问题。

PDL：同一流内的部分重叠

最后一个例外。Hopper／B200 支持的 Programmatic Dependent Launch（PDL，可编程依赖启动），在显式配置启动属性、并在 kernel 内正确设置依赖同步之后，可以让同一 Stream 里前后两个 kernel 部分重叠。它从计算能力 9.0 开始提供，普通 <<< >>> 不会自动启用，也不保证一定产生重叠。PDL 官方说明

四、回到那个向量加法

n = 1000，回到最早那次启动：

CPU 启动一个 Grid，其中有 4 个 Block，也就是 4 个 CTA。
每个 CTA 有 256 个线程，对应 8 个 Warp；整个 Grid 共 32 个 Warp。
硬件根据可用资源，把 CTA 分配到 SM，允许分批执行。
SM 内的 Warp Scheduler 选择就绪 Warp，发射计算或访存指令。
Block 3 的最后一个 Warp 对应元素 992～1023，只有前 8 个线程通过边界检查，完成实际的数组加法。

整个例子没有用到共享内存、Cluster 或 CLC，但把组织方式和执行资源的关系走完了。

概念	回答的问题	与其他概念的关系
Grid	这次启动有哪些线程块？	包含所有 Block，可按 Cluster 分组
Cluster	哪些 Block 需要额外协作？	包含多个 CTA，协同调度在一个 GPC 内
Block／CTA	哪些线程可以在块内协作？	包含多个 Warp，由一个 SM 执行
Warp	哪些线程共同参与指令执行？	同一 CTA 内按 32 个线程分组
Thread／Lane	谁处理具体数据？	Thread 是线程实例，Lane 是其 Warp 内位置
SM	任务在哪里执行？	承载 CTA，调度 Warp，提供执行资源
CLC	运行中的任务怎样领取后续工作？	取消尚未启动的工作，并返回其索引
Stream	多次启动之间怎样排序？	提交 Grid，用 Event 建立跨流依赖

这张表里的每个问题都不同，答案不能互相替代。少了任何一层，另一层的行为就解释不通：只看 Grid 和 Block，无法说明为什么线程多了反而更慢；只看 Warp，也无法说明多个 kernel 为什么排成了队。

五、在 B200 上

B200 的计算能力是 10.0，对应 SM100。线程组织方式没有变：Warp 还是 32 个线程，一个 CTA 还是由一个 SM 执行。变的是矩阵指令、数据搬运和任务调度能力。NVIDIA 计算能力列表

能力	在 B200 上如何使用	适合的工作
Cluster／分布式共享内存	声明 Cluster，使用组内同步和远端共享内存访问	跨 CTA 交换中间数据
TMA	用张量描述符和异步流水线搬运数据	规则的多维张量分块
`tcgen05`／Tensor Memory	使用 SM100 Tensor Core kernel	GEMM、Attention 等矩阵计算
2-SM MMA	选择双 CTA 协作的矩阵指令与调度策略	合适形状的大矩阵分块
FP4／FP8 与块缩放	准备低精度数据及缩放因子，选择对应算子	可接受量化误差的训练或推理
CLC	让持久化 CTA／Cluster 动态领取工作	大量 tile、工作量不均衡

Cluster 和 TMA 是从 Hopper 继承的；tcgen05、Tensor Memory、2-SM MMA、CLC 以及新的低精度矩阵能力，是 Blackwell 扩出来的部分。下面只谈单 GPU 的内核编程，多卡互联不改变这些线程层级。

先选对编译目标

CUDA 12.8 开始能为 SM100 生成原生代码。普通 CUDA kernel 用 -arch=sm_100；要用 SM100 专属指令的 kernel 通常得用 -arch=sm_100a，后面的 CUTLASS 示例就是如此。库可能还要求更高版本的 Toolkit，以它自己的构建要求为准。Blackwell 兼容性指南

# 确认实际使用的 Toolkit；nvidia-smi 显示的是驱动支持的 CUDA 版本上限
nvcc --version

# 假设前面的向量加法代码保存为 add.cu，编译为目标文件
nvcc -O3 -arch=sm_100 -c add.cu -o add.o

100a 里的 a 表示开放架构专属能力，不是“更高一档的优化”。这种目标不提供普通 PTX 那样的跨架构前向兼容保证，Hopper 的 sm_90a 代码也不能拿来替代 B200 的实现。B200 的 SM100 与消费级 Blackwell 的 SM120 同样要区分开。

用两块 CTA 试一次 Cluster 协作

下面每两个 Block 组成一个 Cluster。每块先把自己的一个输入值放进共享内存，再读同组另一块的值，两块都输出这对输入的和。

#include <cuda_runtime.h>
#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void __cluster_dims__(2, 1, 1)
pair_sum(const float* in, float* out) {
    cg::cluster_group cluster = cg::this_cluster();
    __shared__ float local;

    if (threadIdx.x == 0) local = in[blockIdx.x];
    cluster.sync();  // 确保两块都已初始化共享内存

    if (threadIdx.x == 0) {
        int peer_rank = static_cast<int>(cluster.block_rank()) ^ 1;
        float* peer = cluster.map_shared_rank(&local, peer_rank);
        out[blockIdx.x] = local + *peer;
    }

    cluster.sync();  // 远端读取完成前，任何一块都不能退出
}

// d_in、d_out 各有 8 个元素；此例固定启动 8 个 Block
void launch_pair_sum(const float* d_in, float* d_out) {
    pair_sum<<<8, 128>>>(d_in, d_out);
}

输入是 [0,1,2,3,4,5,6,7] 时，预期输出是 [1,1,5,5,9,9,13,13]。存成 cluster_pair.cu 后可以这样编译：nvcc -O3 -std=c++17 -arch=sm_100 -c cluster_pair.cu。显存分配、拷贝、同步和错误检查仍由主程序负责。这个例子每块只处理一个数，用来看协作关系，不适合比较性能。

__cluster_dims__(2,1,1) 固定了每个 Cluster 的块数，Grid 维度必须与之匹配；这里 8 个 Block 形成 4 个 Cluster。block_rank() 是块在 Cluster 内的编号，和 Grid 内的 blockIdx 不是一回事。两次 cluster.sync() 分别守住“数据准备好”和“远端读取结束”。Cluster 编程与分布式共享内存

B200 每 SM 有 228 KB 共享内存，单个 Block 最多可寻址 227 KB；动态共享内存超过 48 KB 要显式 opt-in，用 cudaFuncSetAttribute 设置 cudaFuncAttributeMaxDynamicSharedMemorySize，启动时再传入字节数。Cluster 越大、共享内存越多，能同时驻留的数量越受限，可以用 cudaOccupancyMaxActiveClusters 先估一下配置。Blackwell 调优指南

TMA、Tensor Memory 与 2-SM MMA

TMA（Tensor Memory Accelerator，张量内存加速器）负责搬运，Tensor Core 负责矩阵计算，TMEM（Tensor Memory，张量内存）保存这类矩阵计算的累加结果。TMA 和 TMEM 名字相近，用途不同。

TMA 从 Hopper 开始提供张量异步搬运。用的时候先准备描述形状、步幅和布局的描述符，再发起请求，靠完成屏障判断数据到没到。少量线程就能发起一大块传输，其他 Warp 可以继续处理已经到达的数据。TMA 使用说明

Blackwell 的 tcgen05.mma 把累加结果放进 TMEM。TMEM 是独立的片上存储，需要相应的分配、访问和释放操作，不是声明一个 __shared__ 数组就能拿到，也不属于普通的 HBM 显存分配。第五代 Tensor Core 与 Tensor Memory



%%{init: {"flowchart": {"nodeSpacing": 20, "rankSpacing": 26, "padding": 12}, "themeVariables": {"fontSize": "15px"}}}%%
flowchart TB
    hbm["HBM 显存<br/>矩阵 A、B"]
    smem["共享内存中的输入 tile<br/>完成屏障确认可用"]
    tmem["Tensor Memory<br/>保存矩阵累加结果"]
    epilogue["后处理 Warp<br/>读取结果、缩放或融合计算"]
    output["HBM 显存<br/>输出矩阵 D"]
    hbm -->|TMA 异步搬运| smem
    smem -->|tcgen05.mma 矩阵计算| tmem
    tmem --> epilogue --> output

图 8：一种常见的 B200 GEMM 数据流。省略了多级缓冲、指令描述符和部分访存步骤；不同 kernel 的布局与写回方式可以不同。

这个流程用的就是前面的 Warp 分工：不同 Warp 分别搬运、发起矩阵运算、处理结果。TMEM 让累加结果不必一直占着计算线程的通用寄存器，也方便把“发起矩阵运算”和“处理结果”拆给不同 Warp；生产者与消费者之间仍要靠流水线同步交接。具体分几个角色、各配几个 Warp 由 kernel 决定，前面“1 个搬运 Warp、3 个计算 Warp”不是 B200 的固定配置。SM100 官方示例对 TMEM 与 Warp 分工的说明

2-SM MMA 让同一 Cluster 里的一对 CTA 在两个 SM 上合作完成一次矩阵分块运算。每个 CTA 仍留在自己的 SM 上，跨 SM 的是协作。声明 Cluster 只提供了组织方式，要用双 SM 的矩阵运算，还得选对应的 tcgen05 指令、合法的 tile 与数据布局和同步方式；在 CUTLASS 里对应 1-SM／2-SM 两种调度策略。单纯加大 blockDim 不会打开它。SM100 GEMM 配置

从官方 GEMM 示例入手

上面这些流水线手写起来细节很多。更容易复现的入口是 CUTLASS 的 SM100 示例，下面固定 CUTLASS v3.8.0 和 CUDA 12.8 或更新的环境，避免示例名跟着分支变化；运行需要 B200 和匹配的 NVIDIA 驱动。

git clone --branch v3.8.0 --depth 1 https://github.com/NVIDIA/cutlass.git
cd cutlass
cmake -S . -B build-b200 \
  -DCUTLASS_NVCC_ARCHS=100a \
  -DCUTLASS_ENABLE_TESTS=OFF
cmake --build build-b200 --target 70_blackwell_fp16_gemm -j 4
./build-b200/examples/70_blackwell_gemm/70_blackwell_fp16_gemm \
  --m=4096 --n=4096 --k=4096

这个 FP16 示例用到了新的矩阵指令、TMEM、Warp 分工，以及基于 CLC 的动态调度。跑起来先看结果校验输出 Disposition: Passed，再看耗时。示例的运行时间用来理解实现、比较自己的配置，不宜当作 B200 的峰值性能。官方 FP16 示例与版本化源码

读源码可以顺三条线：ClusterShape_MNK 决定 Cluster 的 CTA 形状，主循环策略决定矩阵计算和搬运怎么组织，tile scheduler 决定下一份工作从哪来。这个版本的 GemmUniversal 默认使用基于 CLC 的 tile scheduler：示例把第四个模板参数写成 void，源码注释注明它默认走 ClusterLaunchControl。普通 CUDA kernel 不会因为编译目标换成 SM100 就自动获得这一领取工作的循环。

FP4／FP8：数据与缩放因子一起准备

B200 支持多种低精度矩阵运算，其中块缩放把一组数值和一份缩放因子绑在一起表达。这里的“块”是量化数据的分组，与 CUDA 的 Block／CTA 没有对应关系。以 NVFP4 为例，算子要的是打包好的 FP4 数据和配套的缩放因子布局，把 FP16 指针换个 4-bit 类型并不会得到正确结果。

FP8 的入口是同目录下的 70_blackwell_fp8_gemm 示例，用同样的方式指定构建目标。NVFP4 用单独的块缩放示例，在同一个 CUTLASS 工作目录里继续：

1
2
3

cmake --build build-b200 --target 72b_blackwell_nvfp4_nvfp4_gemm -j 4
./build-b200/examples/72_blackwell_narrow_precision_gemm/72b_blackwell_nvfp4_nvfp4_gemm \
  --m=4096 --n=4096 --k=4096

这个示例展示低精度输入、缩放因子和矩阵指令怎样配合。接到真实模型还需要量化和校准流程，并检查任务精度；示例的数值校验替代不了模型质量评估。FP8、NVFP4、MXFP8 的编码和缩放规则也不能混用。官方 NVFP4 示例

回到开头：普通向量处理先用基础 CUDA kernel；需要跨块交换中间数据时再考虑 Cluster；矩阵密集的计算从成熟的 SM100 算子和这些示例入手。继续优化时，再分别判断搬运、矩阵吞吐、后处理和任务分配哪一环是瓶颈——一次把所有功能都打开，通常只是多了管理成本。