高级能力的共同目标

现代 GPU 的难点不只是算得快，而是让数据搬运、同步和计算形成流水线。Thread Block Cluster、Distributed Shared Memory、异步 barrier、TMA、异步矩阵计算等能力，都是在扩大可协作范围或降低协调成本。

传统 block：
global → shared → registers → ALU/Tensor Core → shared/global
            ↑ __syncthreads()，阶段常串行

高级流水线：
stage k+1: global ──async copy/TMA──→ shared buffer B
stage k:                              compute buffer A
                                      ↓
                               barrier 交换 A/B

高级 API 不能修复低算术强度、错误布局或过小问题规模。先完成 08、09 章的测量，再引入复杂机制。

能力与最低架构概览

表格只是方向。具体数据类型、方向、维度、原子和 multicast 支持还受架构与 Toolkit 版本限制，必须查目标 GPU 的 Compute Capabilities/Tuning Guide 并运行 capability query。

Cooperative Groups

Cooperative Groups 将参与者作为对象，避免“这个函数暗中假设整个 block 都会到 barrier”的脆弱接口。

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

__device__ float block_sum(cg::thread_block block, float x) {
    extern __shared__ float smem[];
    smem[block.thread_rank()] = x;
    block.sync();

    for (unsigned s = block.size() / 2; s > 0; s >>= 1) {
        if (block.thread_rank() < s)
            smem[block.thread_rank()] += smem[block.thread_rank() + s];
        block.sync();
    }
    return smem[0];
}

组可以是 block、tiled partition、coalesced group、cluster 或 cooperative grid。同步只覆盖该 group，且所有预期成员必须以合法控制流参与。

Grid-wide 同步

普通 kernel 的不同 block 可能分批驻留；在 kernel 内自旋等待所有 block，可能因未调度 block 无法前进而死锁。Cooperative launch 保证整个 grid 可协同驻留，从而允许 grid.sync()。

__global__ void two_phase(float* x) {
    cg::grid_group grid = cg::this_grid();
    phase1(x);
    grid.sync();
    phase2(x);
}

代价是 grid 大小受同时驻留能力约束。若两个阶段之间状态很少，拆成两个 kernel 往往更简单，kernel 边界天然提供 device work 顺序和全局可见性。

Thread Block Cluster

Compute Capability 9.0 引入 cluster，位于 grid 与 block 之间：

Grid
├─ Cluster 0（保证同驻留于一个 GPC）
│  ├─ Block 0 → SM a ┐
│  ├─ Block 1 → SM b ├─ cluster.sync + Distributed Shared Memory
│  └─ Block 2 → SM c ┘
└─ Cluster 1

可移植的 cluster size 上限需要遵循官方规则，并用 occupancy API 查询具体 GPU/MIG 配置允许的大小。更大的 cluster 会减少可同时调度的 cluster 数量，未必更快。

固定 cluster shape

__global__ __cluster_dims__(2, 1, 1)
void clustered_kernel(float* out) {
    namespace cg = cooperative_groups;
    cg::cluster_group cluster = cg::this_cluster();

    // 每个 block 完成本地初始化后再进行跨 block 访问。
    cluster.sync();
    // ... DSM access ...
    cluster.sync();
}

clustered_kernel<<<grid, block>>>(out);

gridDim 仍以 block 数表示。grid 在每个维度上要与 cluster shape 的组织要求匹配。

运行时 cluster shape

cudaLaunchConfig_t cfg{};
cfg.gridDim = grid;
cfg.blockDim = block;

cudaLaunchAttribute attr{};
attr.id = cudaLaunchAttributeClusterDimension;
attr.val.clusterDim.x = 2;
attr.val.clusterDim.y = 1;
attr.val.clusterDim.z = 1;
cfg.attrs = &attr;
cfg.numAttrs = 1;

CUDA_CHECK(cudaLaunchKernelEx(&cfg, clustered_kernel, out));

扩展 launch API 还能承载其他属性。不要同时用冲突的编译期和运行时 cluster 配置。

Distributed Shared Memory

Cluster 内每个 block 仍拥有自己的 shared memory，但其他 block 可把某个 shared pointer 映射到目标 block 的 shared segment：

Block 0 shared ─┐
Block 1 shared ─┼─ 逻辑 DSM address space
Block 2 shared ─┘

任一 cluster block可：read / write / atomic 到成员 block 的 shared memory

典型用例是直方图：每个 block 的局部 bins 不够容纳整个范围时，把 bins 分散到 cluster 多个 block，而不是直接争用 global atomics。

生命周期规则非常关键：

1. 目标 block 必须仍然存在，才能访问其 shared memory；
2. 首次跨 block 使用前，cluster 成员应完成初始化并同步；
3. 任何 block 退出前，所有 DSM 访问都必须完成；
4. DSM 容量、bank 和远端延迟仍需 profile。

异步编程对象

CUDA 的异步 barrier/pipeline 把“发起事务”和“等待事务”分开：

producer                 barrier/pipeline                consumer
reserve stage ───────────────┐
issue async copy             │ expected transaction bytes
do independent work          │
                             ├── completion arrives
wait/acquire stage ◀─────────┘
consume shared data
release stage

常用抽象包括 cuda::barrier、cuda::pipeline、Cooperative Groups copy，以及更底层的 cuda::ptx 包装。选择最高层且能表达所需语义的接口，只有在 profiler 证明必要时再下沉。

cuda::memcpy_async 双缓冲骨架

#include <cuda/pipeline>

template<int TILE>
__global__ void pipelined(const float* in, float* out, int n) {
    __shared__ float buffer[2][TILE];
    auto block = cooperative_groups::this_thread_block();
    auto pipe = cuda::make_pipeline();

    // 示意：真实代码需处理每阶段边界、对齐和参与线程。
    for (int tile = blockIdx.x; tile * TILE < n; tile += gridDim.x) {
        int stage = tile & 1;
        pipe.producer_acquire();
        cuda::memcpy_async(block, buffer[stage],
                           in + tile * TILE,
                           cuda::aligned_size_t<16>(TILE * sizeof(float)),
                           pipe);
        pipe.producer_commit();

        pipe.consumer_wait();
        compute_tile(buffer[stage], out, tile, n);
        pipe.consumer_release();
    }
}

这是结构示意，不是可直接套用的万能循环。真实多阶段 pipeline 要先填充、进入稳态、最后排空；尾部不足对齐大小时需要单独处理，不能伪造 aligned_size_t 承诺。

TMA：描述符驱动的多维 bulk copy

传统线程合作搬运需要每个线程计算地址。TMA 用 tensor map 描述 base、shape、stride、box 和 swizzle 等信息，由少量线程发起硬件批量搬运：

Host 创建 tensor map descriptor
    base = global tensor
    rank/shape/stride
    box dimensions
    element type / interleave / swizzle
                │
Kernel 中一个 elected thread 发起 TMA
                ▼
global tensor tile ──TMA──→ shared / DSM
                │ transaction barrier
                ▼
整个 block/cluster 消费

适合规则的大 tile 和多维布局。对少量、散乱或每线程独特的访问，普通 load/LDGSTS 可能更合适。

TMA 正确性要点

• descriptor 的 global base、shape 与 stride 必须匹配真实 allocation；
• global/shared 地址、barrier 地址和 transfer size 有明确对齐要求；
• 推荐由 elected thread 或 cooperative_groups::invoke_one 发起，避免编译器无法证明单发起者而产生串行；
• global→shared 的完成通常由 transaction barrier 跟踪；
• shared→global 的完成/可复用语义与读取方向不同，要等待正确的 async group；
• shared 写入与 async proxy 之间可能需要 proxy fence；
• 尾块、负 stride、不规则布局不能凭直觉套 descriptor。

Multicast 与 Cluster

支持的平台上，TMA multicast 可将一个 global tile 分发到 cluster 内多个 block 的 shared memory，适合多个 consumer 复用同一数据：

                   ┌→ Block 0 shared → compute 0
global tile → TMA ─┼→ Block 1 shared → compute 1
 multicast         └→ Block 2 shared → compute 2

它减少重复 global 流量，但要求 cluster、目标 mask、shared 布局和 barrier 配合。先确认重复读取确实是瓶颈，再承担额外复杂性。

Programmatic Dependent Launch

普通同 stream kernel B 必须等 A 完全结束才启动。Programmatic Dependent Launch 允许 A 在满足条件后触发 B 的前置、与输入无关的工作，但 B 真正消费 A 结果前仍要执行对应等待。

Kernel A: [setup][produce main result][tail]
                         │ trigger
Kernel B:                [independent setup][wait][consume result]

它不是忽略依赖的并发开关。若 B 在 wait 前错误读取 A 的输出，就是竞态；而且收益取决于可重叠前置工作是否足够大。

Dynamic Parallelism

GPU kernel 可从 device 侧 launch 子 kernel，适合递归细化、数据相关工作生成等场景。但 device launch 具有控制和资源成本，复杂依赖也更难调试。

优先比较：

• host 端批量 launch 或 CUDA Graph；
• persistent kernel + work queue；
• device graph launch；
• dynamic parallelism。

不要用大量极小 child kernel 模拟普通函数调用。

Persistent Kernel 与 Work Stealing

Persistent kernel 启动与 SM 数量相近的 blocks，从全局队列反复取任务：

global work queue: [task0][task1][task2]...[taskN]
                      ↑      ↑
persistent blocks:  fetch → process → fetch → process

适合任务长度不均、需要低延迟调度或避免海量 launch 的工作。难点包括公平性、队列原子热点、退出协议、抢占和与其他 workload 共存。官方 work stealing 能力可减少部分调度开销，但仍必须根据目标架构查询支持。

L2 持久化窗口

对反复访问的小型 global 区域，可以通过 access policy window 提示 L2 保留倾向：

L2 capacity
├─ persisting region：热点查找表/权重
└─ streaming region：一次性数据

这只是缓存策略提示。窗口过大、多个 stream 竞争或工作集超过 set-aside 时，收益会下降。测量实际 L2 hit rate 和 kernel 时间，并在阶段切换时重置策略。

Memory Synchronization Domains

某些大型系统中，一个 kernel 的 system-scope memory fence 会等待与自身无关的远端/NVLink 写入完成，形成 fence interference。Memory synchronization domains 可把相关流量分组，减少不必要的跨域等待。

它是高级系统优化，不改变算法的 happens-before 要求。错误划分可能破坏可见性；只有当 profiler 与通信模式明确指向 fence interference 时才使用。

选型矩阵

调试与验证

1. 先写同步、单缓冲、无 cluster 的正确基线。
2. 加 capability query，对不支持设备保留 fallback。
3. 每次只加入一个机制：async copy、双缓冲、cluster、TMA。
4. 用 Compute Sanitizer 的 race/sync 检查器覆盖边界与部分 tile。
5. 用 Nsight Compute 验证异步拷贝、barrier stall、shared/DSM 和 Tensor 指令。
6. 用 Nsight Systems 检查 kernel 间 overlap 是否真的出现。

练习

1. 把同步 global→shared stencil 改为两阶段 memcpy_async pipeline，绘制 fill/steady/drain 时间图。
2. 用 cluster DSM 实现分片 histogram，与 global atomic 和单 block shared 版本比较。
3. 对二维 copy 比较线程协作 load 与 TMA；改变 tile shape、stride 和尾块。
4. 构造 PDL 的前置独立工作，验证 wait 放错位置时的错误，再测量正确版本的 overlap。

官方参考

• CUDA Programming Guide: Advanced Kernel Programming
• CUDA Programming Guide: Programming Model / Thread Block Clusters
• CUDA Programming Guide: Asynchronous Data Copies
• CUDA Programming Guide: Programmatic Dependent Launch
• CUDA C++ Core Libraries: Synchronization
• Hopper Tuning Guide
• Blackwell Tuning Guide