Cuda

CUDA C++ 编程精要

按执行模型、内存体系、异步执行、特性专题与性能准则梳理 CUDA C++ 的核心概念，附术语表与官方文档章节索引。

本文是对《CUDA Programming Guide》的阅读整理。

1. 阅读路径

CUDA 的知识点彼此关联：不懂 SIMT 就理解不了分支分歧的代价，不懂内存层次就理解不了合并访问，不懂内存模型就写不对跨块同步。因此建议按下面的依赖顺序阅读，而不是按 API 字母序查。

2. 异构执行模型

2.1 为什么是 CPU + GPU

CPU 与 GPU 的设计目标不同，芯片资源分配因此分道扬镳 [1][2]：

维度CPUGPU
设计目标让单条线程尽量快让数千条线程同时在跑（吞吐优先）
并行规模数十条线程数千条线程
晶体管去向大缓存与复杂流控数据处理单元（浮点等）
隐藏延迟的手段缓存 + 流控用海量并行掩盖访存延迟

GPU 用更多晶体管"算"，而不是"等"。但真实应用总是并行段与串行段的混合，所以系统形态是 CPU + GPU 各司其职：CPU 负责控制流与串行逻辑，GPU 负责可并行的密集计算。

2.2 硬件层次：SM 与 GPC

  • GPU 由若干流多处理器（SM）组成，SM 再组成图形处理簇（GPC）。
  • 每个 SM 内含寄存器堆、统一数据缓存、若干功能单元；统一数据缓存同时为共享内存和 L1 缓存提供物理资源，两者的划分比例可以运行期配置。
  • 各架构的 SM 规格（内存大小、功能单元数量）不同，但不影响编程模型的正确性——模型与硬件解耦。

2.3 网格、线程块与线程

一次 kernel 启动会启动成千上万条线程，它们的组织关系是三层结构 [2]：

由此导出两条编程铁律：

  1. 块内可协作：块内线程同驻一个 SM，可通过片上共享内存高效通信，并用 __syncthreads() 同步。
  2. 块间无顺序保证：块的调度顺序（并发还是串行）不受保证，因此块与块之间不得有数据依赖，任何块顺序下结果都必须正确。

启动时通过执行配置指定网格与块的维度（<<<...>>>），可选参数还有动态共享内存字节数与流；簇大小、SM 配置等更复杂的项需用 cudaLaunchKernelEx 传递启动属性 [3][9]。

2.4 warp 与 SIMT

  • 块内线程以 32 条为一组构成 warp，warp 是硬件调度单位，以 SIMT（单指令多线程）方式执行。
  • warp 内每条线程有自己的指令计数器与寄存器状态，因此可以走不同分支；但硬件每次只能发射一条公共指令，于是不同路径被串行执行，其余线程被掩蔽——这就是分支分歧（divergence）。
  • 推论：块内线程数最好是 32 的倍数，否则最后一个 warp 存在空闲 lane，浪费功能单元与访存带宽。
  • SIMT 与 SIMD 的区别：SIMD 只有单一控制流且数据宽度固定；SIMT 允许每线程独立路径，无固定数据宽度 [2]。

2.5 线程块簇（可选分组）

Compute Capability 9.0（Hopper）及以上支持把若干相邻块组成簇 [2][15]：

  • 指定簇不改变网格维度与块索引；簇内所有块同驻一个 GPC，并被一起调度。
  • 簇内块可以互相访问彼此的共享内存，形成分布式共享内存（DSMEM），并用 Cooperative Groups 提供的接口做跨块同步。
  • 簇的可移植最大尺寸为 8 个块，实际上限由硬件决定。

2.6 计算能力（CC）

CC 是设备能力版本号，决定了块内最大线程数、共享内存上限、寄存器数量以及是否支持簇、TMA 等特性 [26]。编译期用虚拟架构与真实架构两个维度控制产物形态，运行期可用 cudaGetDeviceProperties 查询。

架构CC关键新增
Volta7.0独立线程调度、Tensor Core
Turing7.5光线追踪核心
Ampere8.0cp.async 异步拷贝、异步屏障
Hopper9.0TMA、线程块簇、DPX
Blackwell10.x / 12.x新一代 Tensor Core 与新数据类型

3. Kernel、索引与边界

3.1 定义与启动

用 __global__ 声明 kernel，返回类型必须为 void；它由许多线程并行执行 [3]。启动时给出执行配置，前两个参数是网格维度与块维度，一维场景可以直接用整数：

__global__ void vecAdd(float* A, float* B, float* C) { }

// 1 个块、256 条线程
vecAdd<<<1, 256>>>(A, B, C);

多维场景用 dim3；完整形式还包含动态共享内存字节数与流：

kernel<<<gridDim, blockDim, sharedMemBytes, stream>>>(args...);

两点常被忽略：

  • 每块线程数上限为 1024（资源允许时一个 SM 可同时驻留多个块）。
  • kernel 启动相对 host 是异步的：host 不会等它开始或完成就继续执行，需要结果时必须显式同步 [3][6]。

3.2 索引计算

kernel 内通过内置变量确定自己的位置 [3][4]：

变量含义值域
threadIdx块内线程索引（.x/.y/.z）0 .. blockDim-1
blockDim块维度启动时指定
blockIdx块在网格中的索引0 .. gridDim-1
gridDim网格维度启动时指定

一维并行最常见的线性化方式：

int workIndex = threadIdx.x + blockDim.x * blockIdx.x;
C[workIndex] = A[workIndex] + B[workIndex];

多维块与网格只是方便表达，本身不影响性能；线程线性化规则是 x 变化最快、y 步长为 blockDim.x、z 步长为 blockDim.x * blockDim.y，这条规则决定了线程如何被划分进 warp [4]。

3.3 边界检查与块数计算

当数据长度不是块大小的整数倍时，必须做边界检查，并把块数上取整：

__global__ void vecAddChecked(float* A, float* B, float* C, int n) {
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

int threads = 256;
int blocks  = (n + threads - 1) / threads;   // 上取整
vecAddChecked<<<blocks, threads>>>(devA, devB, devC, n);

CCCL 提供了 cuda::ceil_div（头文件 <cuda/cmath>）替代手写上取整 [32]。多启动几个"不做事的线程"代价很小，但应避免启动整块都无工作的块；256 线程/块是一个稳妥的起点。

3.4 一个完整可运行的示例

把前几节拼起来，就是下面这个能编译、能自校验的最小程序。语法上它就是 C++，但每一处与标准 C++ 不同的地方，都在注释里标了出来——差异集中在三点：执行空间、执行配置 <<<...>>>、以及指针属于哪一侧的内存。

// CUDA C++ 的"非标准"之处：这段代码有三处是标准 C++ 里不存在的写法——
// 执行空间说明符（__global__）、执行配置（<<<...>>>）、设备内存 API（cuda*）。
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

// CUDA 的 API 是 C 风格接口：返回错误码，不像标准 C++ 那样抛异常，
// 所以每个调用都要自己检查返回值，否则错误会一路静默传递。
#define CUDA_CHECK(call)                                                    \
    do {                                                                    \
        cudaError_t err = (call);                                           \
        if (err != cudaSuccess) {                                           \
            std::fprintf(stderr, "%s:%d %s\n", __FILE__, __LINE__,          \
                         cudaGetErrorString(err));                          \
            std::exit(EXIT_FAILURE);                                        \
        }                                                                   \
    } while (0)

// __global__ 是执行空间说明符：标准 C++ 没有这个概念。
// 它表示"编译为 GPU 代码、由 host 启动、在 device 上执行"，且必须返回 void。
__global__ void vecAdd(const float* A, const float* B, float* C, int n) {
    // threadIdx / blockDim / blockIdx 不是标准库变量，而是编译器内置变量：
    // 语法与普通变量一样，但每条线程各有一份，用它算出自己的唯一编号。
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    // 块大小通常不能整除数据量，最后一块里的部分线程会越界，必须挡住。
    if (i < n) {
        C[i] = A[i] + B[i];     // 这里的指针指向显存，解引用由 GPU 完成
    }
}

int main() {
    // 从这里往下是 host 代码，由 CPU 执行，写法与普通 C++ 一样。
    // 但 host 与 device 的内存是分开的：malloc / new 只对 host 侧有效。
    const int    n     = 1 << 20;                 // 1M 个元素
    const size_t bytes = n * sizeof(float);

    float* hA = static_cast<float*>(std::malloc(bytes));
    float* hB = static_cast<float*>(std::malloc(bytes));
    float* hC = static_cast<float*>(std::malloc(bytes));
    for (int i = 0; i < n; ++i) { hA[i] = 1.0f; hB[i] = 2.0f; }

    // 为什么非要手写这些拷贝？因为 GPU 有自己的显存，与 host 的系统内存是
    // 两块物理内存：kernel 跑在 device 侧，只能访问设备内存。
    // 把 host 指针当普通 C++ 指针用（在 device 上解引用）不是"慢一点"，
    // 而是非法访问。所以流程必然是：分配设备内存 -> 拷进去 -> 启动 -> 拷回来。
    float *dA = nullptr, *dB = nullptr, *dC = nullptr;
    CUDA_CHECK(cudaMalloc(&dA, bytes));
    CUDA_CHECK(cudaMalloc(&dB, bytes));
    CUDA_CHECK(cudaMalloc(&dC, bytes));
    // cudaMemcpy 的参数是 C 风格 void*，看不出方向，所以要显式指定
    // cudaMemcpyHostToDevice；若系统支持统一虚拟地址空间，也可传
    // cudaMemcpyDefault 让驱动按指针推断（见第 5 节）。
    CUDA_CHECK(cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(dB, hB, bytes, cudaMemcpyHostToDevice));

    // <<<blocks, threads>>> 不是 C++ 语法，而是 nvcc 的执行配置扩展：
    // 标准 C++ 里没有"一次调用启动 N 个线程"的表达方式。
    const int threads = 256;
    const int blocks  = (n + threads - 1) / threads;
    vecAdd<<<blocks, threads>>>(dA, dB, dC, n);

    // 关键差异：启动是异步的。这一行返回时 kernel 可能还没开始执行，
    // 与标准 C++"函数返回即执行完毕"的直觉相反，因此错误分为两处查。
    CUDA_CHECK(cudaGetLastError());                // 捕获启动期错误
    CUDA_CHECK(cudaDeviceSynchronize());           // 等待执行完成

    CUDA_CHECK(cudaMemcpy(hC, dC, bytes, cudaMemcpyDeviceToHost));
    bool ok = true;
    for (int i = 0; i < n && ok; ++i) { ok = (hC[i] == 3.0f); }
    std::printf("result: %s\n", ok ? "PASS" : "FAIL");

    // 设备内存必须显式释放；host 内存照旧走标准 C++ 的 free。
    CUDA_CHECK(cudaFree(dA));
    CUDA_CHECK(cudaFree(dB));
    CUDA_CHECK(cudaFree(dC));
    std::free(hA);
    std::free(hB);
    std::free(hC);
    return 0;
}

nvcc -O2 -arch=native vector_add.cu -o vector_add
./vector_add

为什么这段代码必须手写两次 cudaMemcpy：CUDA 不只是"给 C++ 加几个关键字"，它把一台机器的内存劈成了两半，而两半之间不会自动搬运数据。

host 侧（CPU + 系统内存）                 device 侧（GPU + 显存）
  hA / hB / hC                              dA / dB / dC
  由 malloc 分配                           由 cudaMalloc 分配
       |                                          ^
       +---- cudaMemcpy(..., HostToDevice) -------+
       +---- cudaMemcpy(..., DeviceToHost) -------+
  • kernel 在 device 上执行，只能访问设备内存。把 host 指针传进 kernel 再解引用，不是"慢一点"，而是非法访问（报 invalid memory access）。
  • 所以"数据先在 GPU 能访问的内存里"是硬约束，不是风格选择。标准 C++ 里的引用、指针、std::copy 全部只在同一地址空间内成立，跨不过这条边界。
  • cudaMemcpy 的第 4 个参数正是为这件事服务的：接口是 C 风格的 void*，看不出方向，所以要显式写 cudaMemcpyHostToDevice 或 cudaMemcpyDeviceToHost。系统支持统一虚拟地址空间时，也可以传 cudaMemcpyDefault 让驱动按指针自动推断（见第 5 节）。
  • 想彻底省掉这份手写拷贝，可以改用托管内存（cudaMallocManaged），由驱动按需迁移；代价是数据放置不可控、性能更依赖系统能力。两种路线的取舍在第 5 节有对照表。

四点值得留意：

  • cudaGetLastError() 要紧跟启动调用，它把"启动失败"（比如块尺寸越界）与"执行中出错"区分开；漏掉它，错误会漂移到后面某次同步调用上，排查时会指向错误的位置。
  • 校验循环只是示意。真实工程应比较浮点误差，而不是要求逐位相等。
  • -arch=native 按本机 GPU 生成 Cubin；若要分发给别的机器，应改为生成 PTX 或显式列出多个目标（见第 9 节）。
  • 中文 Windows 上还有一类经典陷阱：源文件是 UTF-8 而系统代码页是 GBK 时，MSVC 会把中文注释末尾的某些字节当成续行符，于是注释吞掉下一行代码，报出 identifier "bytes" is undefined 这种看起来无缘无故的错误。加 -Xcompiler /utf-8 即可（也可以把源文件保存为带 BOM 的 UTF-8）：
nvcc -O2 -arch=native -Xcompiler /utf-8 vector_add.cu -o vector_add

4. 内存体系

4.1 层次、作用域与生命周期

内存类型作用域生命周期物理位置
寄存器线程线程SM
本地内存线程线程显存（逻辑私有，物理在全局）
共享内存块kernelSM 片上（与 L1 共享资源）
全局内存网格 + host应用显存
常量内存网格（只读）应用显存（带缓存）

访问代价按"距离"递增：寄存器最快，共享内存低延迟，全局内存容量大但延迟高 [2][4]。

落到硬件上，一个 SM 的组织大致是这样：

SM
+-- 寄存器堆                每线程私有，编译器分配
+-- 统一数据缓存            物理资源可配置切分
|     +-- 共享内存          块内共享，用户显式管理
|     +-- L1 缓存           由硬件管理
+-- 功能单元                FP32 / INT / Tensor Core / 访存单元 ...
+-- warp 调度器             每周期挑一个就绪 warp 发射指令

共享内存与 L1 抢的是同一块物理资源，这既是"多给共享内存能加速"的原因，也是"给过头反而变慢"的原因：L1 被挤小，命中率下降会把压力转移回全局内存。

4.2 共享内存：静态与动态

__shared__ float staticArray[1024];             // 静态：编译期定长

kernel<<<grid, block, sharedBytes>>>();          // 启动时给出字节数
extern __shared__ float dynamicArray[];          // 动态：一个 kernel 只能有一个

动态共享内存每个 kernel 只能有一个；需要多块视图时，申请一个大数组自行切分，并注意对齐：

extern __shared__ float array[];
short* a0 = (short*)array;            // 2 字节对齐
float* a1 = (float*)&a0[128];         // 128*2 = 256 字节，满足 4 字节对齐
int*   a2 = (int*)&a1[64];

若切分位置未按目标类型对齐（例如把 float* 落在 254 字节偏移处），访问行为不正确 [4]。

4.3 两种访存效率陷阱

未合并的全局访问：同一 warp 的 32 条线程若访问连续地址，硬件会把它们合并成尽可能少的内存事务；地址离散则事务数暴涨、吞吐骤降。设计 kernel 时应让 threadIdx.x 相邻的线程访问相邻地址 [4]：

float v = data[threadIdx.x + blockIdx.x * blockDim.x];   // 合并访问

结构体数组（AoS）是这件事上最常见的坑：相邻线程访问 p[i].x 时地址间隔是 sizeof(Particle)，一个 128 字节的事务里只用得上一小部分；把字段拆成各自连续的数组（SoA）后，同一个 warp 的访问重新变得连续：

struct Particle { float x, y, z; };
Particle* p;            // AoS：相邻线程的地址间隔 sizeof(Particle)
float  x  = p[i].x;     // 事务利用率低

float *px, *py, *pz;    // SoA：每个字段各自连续
float  x2 = px[i];      // 合并访问

共享内存 bank 冲突：共享内存按 bank 组织（当前架构 32 个 bank，相邻 4 字节映射到相邻 bank）。warp 内多条线程访问同一 bank 的不同地址时访问被串行化；所有线程访问不同 bank、或访问同一地址（广播）都不会冲突。数据重排（如先读入共享内存再转置）常用于把不可合并的访问转成可合并访问 [4]。

4.4 寄存器、本地内存与占用率

  • 寄存器由编译器分配，可用 -maxrregcount 或 __launch_bounds__ 限制每线程上限；限制过紧会导致寄存器溢出（spilling），值被写到全局内存再读回。
  • 本地内存是逻辑私有、物理位于全局内存的存储，常用于索引不确定的数组、放不进寄存器的结构体以及溢出变量，访问代价与全局内存同级 [4]。
  • 占用率 = SM 上活跃 warp 数 / SM 支持的最大 warp 数，受每线程寄存器数、每块共享内存量、块与线程数上限共同约束；资源不足以驻留至少一个块时 kernel 直接启动失败。它是"够用即好"的指标，不是越高越好——寄存器换并行度、共享内存换 L1 容量都需要权衡 [4]。

5. 数据放置：统一内存与显式管理

数据位于访问它的处理器本地内存时性能最佳。CUDA 给了两条路：让驱动自动搬迁（统一内存），或自己显式管理（cudaMalloc + cudaMemcpy）[7]。

5.1 统一虚拟地址空间（UVA）

同一 OS 进程内，host 内存与所有 GPU 显存处在同一个虚拟地址空间。因此：

  • cudaPointerGetAttributes() 可以从指针值反查内存位置（CPU 还是哪个 GPU）；
  • cudaMemcpy* 的 cudaMemcpyKind 可传 cudaMemcpyDefault，由指针自动推断拷贝方向 [7]。

5.2 统一内存的四种范式

统一内存的行为取决于 OS、内核版本、GPU 与互联方式。用三个设备属性即可判定当前系统属于哪种范式 [7]：

属性含义
cudaDevAttrConcurrentManagedAccess1 = 完整统一内存支持；0 = 受限支持
cudaDevAttrPageableMemoryAccess1 = 全部系统内存都统一；0 = 仅显式托管分配统一
cudaDevAttrPageableMemoryAccessUsesHostPageTables1 = 硬件一致性；0 = 软件一致性
int concurrentManagedAccess = 0, pageableMemoryAccess = 0, usesHostPageTables = 0;
cudaDeviceGetAttribute(&concurrentManagedAccess, cudaDevAttrConcurrentManagedAccess, dev);
cudaDeviceGetAttribute(&pageableMemoryAccess, cudaDevAttrPageableMemoryAccess, dev);
cudaDeviceGetAttribute(&usesHostPageTables,
                       cudaDevAttrPageableMemoryAccessUsesHostPageTables, dev);
  • 受限支持：托管内存先落在 CPU 内存，以较粗粒度迁移；GPU 活跃期间 CPU 不得访问；不允许超额订阅。
  • 完整支持：托管内存通常分配在首次触碰它的处理器内存中，被其他处理器使用时按页（软件一致性）或缓存行（硬件一致性）迁移，允许超额订阅显存。
  • 硬件一致性（Grace 系列 + NVLink C2C）：驻留在 GPU 的托管分配可被 CPU 直接访问而无需迁移，还支持 host 侧原生原子；但硬件一致性不会让 host 访问纯 GPU 分配（如 cudaMalloc 的结果）[7]。

托管内存的分配方式有 cudaMallocManaged、cudaMallocFromPoolAsync（池类型为托管）以及 __managed__ 全局变量。性能调优的两个抓手是 cudaMemAdvise（声明只读、偏好位置、将被谁访问）与 cudaMemPrefetchAsync（在 kernel 前异步预取，与计算重叠）[7]。

5.3 取舍

维度显式管理统一内存
代码复杂度高（每处拷贝都要写）低（自动迁移）
性能可预测性高（放置可控）取决于系统能力与提示质量
适用场景性能关键、数据流明确快速开发、复杂数据结构、需要超额订阅

显式管理的典型流程是分配、上行拷贝、执行、下行拷贝、释放；异步拷贝指向 host 的缓冲区必须是页锁定内存，否则会退化为同步行为。页锁定内存用 cudaMallocHost / cudaHostAlloc 分配，用 cudaFreeHost 释放，也可用 cudaHostRegister 把已有内存注册为页锁定 [7]。

6. 异步执行

6.1 可以并发的六件事

host 计算、device 计算、H2D 拷贝、D2H 拷贝、设备内拷贝、设备间拷贝，这些都可以彼此重叠 [6]。异步接口的语义是调用立即返回（可能在工作完成、甚至开始之前），因此需要结果时必须同步：

同步风格行为代表接口
阻塞调用阻塞直到完成cudaDeviceSynchronize
非阻塞立即返回并给出状态cudaStreamQuery
回调完成时执行注册的函数cudaLaunchHostFunc

6.2 流与事件

流是一个命令队列：往里追加拷贝与 kernel，按入队顺序执行；不同流之间可以并发。事件是插入流的标记，用来表达依赖与计时 [6]。

cudaStream_t stream;
cudaStreamCreate(&stream);

kernel<<<grid, block, 0, stream>>>(...);                             // 流内启动
cudaMemcpyAsync(dst, src, bytes, cudaMemcpyHostToDevice, stream);    // 流内异步拷贝

cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);
cudaEventRecord(start, stream);
cudaEventRecord(stop, stream);
cudaStreamSynchronize(stream);
float elapsedMs = 0.0f;
cudaEventElapsedTime(&elapsedMs, start, stop);

cudaEventDestroy(start); cudaEventDestroy(stop);
cudaStreamDestroy(stream);

两个易错点：

  • 异步拷贝涉及 CPU 内存时要求页锁定（pinned），否则无法与其他工作重叠 [6]。
  • 默认流与其他流之间存在隐式同步语义；用 cudaStreamNonBlocking 创建的流不与默认流隐式同步 [6]。

6.3 CUDA Graphs

流中每次 kernel 启动都要付出 host 侧设置开销；对短 kernel，这笔开销可能占据端到端时间的显著比例。图把工作流定义一次、实例化一次，之后重复启动的开销极低，并把整体结构交给 CUDA 系统去做跨操作优化 [14]：

  • 结构：节点是操作，边是依赖；节点依赖满足即可被调度。
  • 构建：Graph API 逐节点构图，或用 Stream Capture 捕获一段流操作自动构图。
  • 执行：cudaGraphInstantiate 实例化后 cudaGraphLaunch 反复启动；图优化的主要成本在实例化期。
  • 更新：cudaGraphExecUpdate 可整图或逐节点更新参数与依赖，避免重新实例化。

6.4 程序化依赖启动（PDL）

普通流模型下，后继 kernel 要等前驱完成、host 才知道可以启动，存在 CPU 往返延迟。PDL 让前驱 kernel 的设备代码自己触发后继 kernel 启动，把 kernel 间依赖交给设备侧控制，从而省掉 CPU 往返 [9][16]。

7. 同步与内存模型

7.1 块内屏障

__syncthreads() 阻塞块内所有线程直到全部到达，并保证此前所有内存访问在屏障后对块内线程可见 [4]：

__global__ void blockReduce(const int* in, int* out) {
    __shared__ int shared[128];
    shared[threadIdx.x] = in[blockDim.x * blockIdx.x + threadIdx.x];
    __syncthreads();                     // 到此为止所有写入对块内可见

    if (threadIdx.x == 0) {
        int sum = 0;
        for (int i = 0; i < blockDim.x; ++i) sum += shared[i];
        out[blockIdx.x] = sum;
    }
}

若块内部分线程因分支未执行 __syncthreads()，行为未定义（可能死锁）。更细粒度的选择是 __syncwarp()（warp 内）、cuda::barrier（可被异步操作到达）与 Cooperative Groups 的组同步 [10][15][17]。

7.2 原子操作

对全局或共享内存的原子读-改-写避免数据竞争，常用 atomicAdd/atomicSub、atomicExch、atomicMin/atomicMax、atomicInc/atomicDec、atomicCAS、位运算系列 [4]。注意：同一 warp 内对同一地址的原子操作会被串行化，且执行顺序不保证（每次读-改-写都会发生）。

7.3 线程作用域

内存操作的可见范围与同步范围由线程作用域定义 [10][28]：

作用域可见/可同步范围
cuda::thread_scope_thread仅本线程
cuda::thread_scope_block同块线程
cuda::thread_scope_device同一设备（同一内存同步域内）
cuda::thread_scope_system系统内所有线程（含 CPU 与其他 GPU）

它把"同步距离"的成本显式化：块内同步便宜、系统级同步昂贵，因此标准 C++ 那种"同步成本均匀"的假设在 CUDA 中不成立 [28]。

7.4 一致性保证的层次

同步需求所需手段
同块内线程__syncthreads() 或 cuda::barrier 建立 happens-before
跨块/网格内设备级作用域原子或栅栏（thread_scope_device），或 cooperative launch
跨设备/与 CPU系统级作用域原子或栅栏（thread_scope_system）；跨内存同步域时还需额外的 fence [21]

内存序方面，cuda::atomic 支持 relaxed/consume/acquire/release/acq_rel/seq_cst，默认为 seq_cst [28]。

7.5 出错时的典型症状

同步写错的麻烦在于它通常不报错，只表现为"结果偶尔不对"。按症状定位：

症状常见原因正确写法
结果偶发错误，换个块尺寸又好了读共享内存前漏了块内屏障写入与读取之间加 __syncthreads()
程序偶发挂死屏障落在分支里，部分线程到不了把 __syncthreads() 移出条件分支，或改用 Cooperative Groups 的组同步
单块跑得对，多块结果错依赖了块间执行顺序用设备级原子或栅栏；需要网格级同步时用 cooperative launch
CC 7.0 前后行为不一致依赖了隐式 warp 锁步用 __syncwarp() 显式同步
CPU 读到旧值缺系统级栅栏就跨设备读写用 thread_scope_system 的原子或栅栏

共同点：这几种写法在某一代硬件上可能"刚好能跑"，但它们依赖的是实现细节，不是内存模型给出的保证。

8. 特性专题速览

下表是整个特性面的索引，按"需要知道完整细节时查官方页"的方式使用；每一项在官方指南都有独立章节 [1][13]。

特性一句话动机入口 API / 关键字
Tile 编程模型以块为单位写代码，把并行度交给编译器，更容易用上 TMA 与 Tensor Core__tile_global__ / __tile__ / ct::tile [5]
Cooperative Groups提供安全、面向未来的组同步与 scan/reduce 原语，替代手写脆弱的 warp 内同步this_thread_block() / tiled_partition / cg::reduce [15]
Programmatic Dependent Launch前一 kernel 设备侧触发后一 kernel，消除 CPU 往返cudaTriggerProgrammaticLaunchCompletion [16]
异步屏障屏障可被异步操作"到达"，支撑生产者-消费者流水线cuda::barrier [17]
流水线多阶段异步搬运，stage count 支持预取多批数据cuda::pipeline + cuda::memcpy_async [18]
异步数据拷贝全局到共享的搬运不再经寄存器；TMA 支持多维张量、im2col、Swizzlecp.async（CC 8.0+）、TMA（CC 9.0+）[19]
CUDA Graphs定义一次、重复低开销启动，并可做跨操作优化cudaGraphCreate / Stream Capture [14]
流序内存分配器分配与释放按流排序，规避 CPU 侧分配锁与 use-after-free 竞态cudaMallocAsync / cudaFreeAsync [1]
L2 缓存控制把 L2 的一部分留给频繁复读的数据cudaLimitPersistingL2CacheSize + accessPolicyWindow [20]
内存同步域跨域原子/栅栏需要额外 fence 才能保证全局排序cudaDevAttrMemorySyncDomainCount [21]
虚拟内存管理把物理分配与虚拟映射、访问权限分离，支持稀疏与自定义分配器cuMemCreate/cuMemMap（仅 Driver API）[22]
进程间通信跨进程共享显存与事件cudaIpcGetMemHandle / cudaIpcOpenMemHandle [24]
CUDA 动态并行设备代码自行启动子网格，适合递归与不规则并行设备侧启动（CDP2，CUDA 12.0 起默认）[23]
图形互操作与 OpenGL/D3D/Vulkan 共享资源，渲染与计算互喂cudaGraphicsGLRegisterBuffer 等 [25]
绿色上下文对 GPU 功耗与频率做上下文级控制cudaGreenCtxCreate [1]
惰性加载把 cubin 加载推迟到首次使用，加快启动CUDA_MODULE_LOADING=LAZY [1]
Driver 入口点访问在 runtime 程序里直接取 driver 函数指针cudaGetDriverEntryPoint [1]

其中两条细节值得单独记：异步代理（async proxy）——TMA 这类批量搬运走 async proxy，其与普通访问之间不保证排序，需要 proxy fence；CUDA 13 起，__global__ 函数与设备变量默认内部链接，跨编译单元引用需要显式 extern [10][8]。

9. 编译与工具链

nvcc 是编译器驱动：对外提供统一命令行，对内按阶段调用宿主 C++ 编译器、PTX 汇编器与 ptxas [8][31]。

扩展名内容
.c / .cpp / .cc / .cxx仅 host 代码
.h / .hpp / .hh / .hxxhost / device / 混合
.cuCUDA 源文件（device + host）
.cuhCUDA 头文件

架构控制（默认生成最低架构以最大化兼容性）[8]：

选项效果
-arch=compute_XY只生成该虚拟架构的 PTX（可 JIT 前向兼容）
-arch=sm_XY生成 PTX + Cubin
-arch=native按本机 GPU 生成 Cubin（无 PTX）
-arch=all全部支持架构的 Cubin + 最新 PTX
-gencode=arch=compute_XY,code=sm_XY精确控制多目标

分离编译：默认是整体编译，使用 GPU 代码的编译单元必须自带全部设备代码。加 -rdc=true（别名 -dc）后设备代码可跨编译单元链接，组织更灵活、编译更快、二进制更小，代价是链接期优化机会减少——可用 LTO（-dlto）找回大部分性能 [8]：

nvcc -dc -dlto -arch=sm_100 definition.cu -o definition.o
nvcc -dc -dlto -arch=sm_100 example.cu    -o example.o
nvcc -dlto definition.o example.o -o program

调试与优化常用组合：-g（host 调试信息）、-G（设备调试信息，会抑制优化）、-lineinfo（行号信息，供 compute-sanitizer 与 Nsight 关联）、-res-usage（资源占用报告）、-Xptxas=-warn-spills（寄存器溢出告警）。性能分析交给 Nsight Compute（kernel 级）与 Nsight Systems（系统级时间线）[33][34]。

10. 性能准则

官方优化策略可以归纳成三句话 [27]：

1. 最大化利用率 (Maximize Utilization)
   应用级: 用流/图重叠拷贝与计算
   设备级: 数据常驻显存, 避免反复 H2D/D2H
   SM 级 : 用占用率找到块/线程平衡点
2. 最大化访存吞吐 (Maximize Memory Throughput)
   减少 host-device 传输(最贵)
   全局访问合并(coalesced)
   复用数据放共享内存, 避免重复读全局
3. 最大化指令吞吐 (Maximize Instruction Throughput)
   避免 warp 内分支分歧
   使用内建快速数学函数
   循环展开, 但权衡寄存器压力

设计动机值得点明：GPU 靠"海量并行 + 合并访存 + 指令级流水"掩盖延迟，所以优化目标不是让单线程更快，而是让整个 SM 的并行度、访存效率与指令效率同时最大化。反过来，任何把线程串起来、把访存打散、把分支岔开的重构，都会直接吃掉这三项收益。

先量再调，否则改的大概率不是瓶颈。三个够用的入口：

手段看什么怎么判断
nvcc -res-usage（或 --ptxas-options=-v）每线程寄存器数、共享内存用量、是否溢出出现 spill 说明寄存器不够，先放宽 __launch_bounds__
cudaOccupancyMaxActiveBlocksPerMultiprocessor实际可驻留的块数与理论上限差得多，说明受寄存器或共享内存限制
Nsight Compute [33]计算与访存吞吐占比、warp stall 原因访存吞吐接近峰值而计算占比低，先改访存；反之先改指令

用 Nsight Compute 时别一上来就看全部指标：先判断 kernel 处于"访存受限"还是"计算受限"，再顺着对应的那组指标往下看。

11. 常见误解

误解实际情况
占用率越高越快占用率只是"够用即好"。提高它通常靠降寄存器或降共享内存，代价可能是溢出与重复访存，净效果可能更差
共享内存给得越多越快共享内存与 L1 共用物理资源，给多了会挤小 L1；bank 冲突没解决时，多给也不会变快
cudaMemcpy 一定慢慢的是不可重叠的同步拷贝。配合页锁定内存与多流，拷贝能与计算重叠，此时它不再是瓶颈
块数越多越好块数不足会让 SM 空转，过多则调度与尾部开销上升。覆盖全部数据且能整除，基本就是合适的量
__syncthreads() 能同步整个网格它只作用于当前块。跨块要用设备级原子或栅栏，网格级同步得用 cooperative launch

12. 结语

CUDA 的每条约束都能追溯到硬件：块内线程同驻一个 SM，所以共享内存可用、块内屏障便宜；块与块之间没有调度顺序，所以跨块同步只能靠全局内存或协同启动；同步成本随线程距离上升，所以线程作用域把这笔代价显式写进了代码。那些"在某一代硬件上刚好能跑"的写法——依赖 warp 锁步、反复重读全局内存、让同一 warp 走不同分支——都会在架构换代时暴露问题。

优化的方向也只有一个：让并行度、访存效率、指令效率同时最大化。剩下的功课是把它落到工具链上：选对编译目标与分离编译策略，用流与图组织异步，需要新特性时按第 8 节的特性索引回查官方文档。

12. 术语表

术语含义
SM流多处理器，块执行的物理单元
GPCGPU 处理簇，SM 的集合
warp32 条线程组成的硬件调度单位
SIMT单指令多线程执行模型
lanewarp 中线程的位置（0-31）
divergencewarp 内分支路径不同导致的串行执行
coalesced access同一 warp 访问连续地址，合并为少数内存事务
occupancySM 上活跃 warp 数与理论上限之比
execution configuration启动配置 <<<grid, block, smem, stream>>>
PTX虚拟指令集，与具体 GPU 无关的中间表示
Cubin / Fatbin原生机器码 / 多架构打包容器
CCCompute Capability，设备能力版本号
UVA统一虚拟地址空间
managed memory统一内存托管分配，可自动迁移
pinned memory页锁定主机内存，异步拷贝的前提
streamGPU 命令队列
event流内标记与时间戳
CDPCUDA 动态并行
TMATensor Memory Accelerator
DSMEM分布式共享内存（簇内跨块访问）
thread scope内存操作的可见性与同步范围
async proxy异步操作的内存访问代理

14. 参考

行内引用编号 [n] 对应下列条目。主文档为《CUDA Programming Guide》分章节版，另有经典版《CUDA C++ Programming Guide》保留传统章节编号，两者内容互补。

链接与章节编号依据 2026 年 9 月的官方文档核对。文中带版本前提的结论（例如 CUDA 13 起设备符号默认内部链接、CDP2 的适用条件）以该版本为准，跨版本使用时请对照官方的变更说明。

© 2026 Sanmusan33 的博客