第 2 章 · Hello CUDA

学习目标

• 写出第一个 __global__ kernel，理解 host 与 device 的分工
• 掌握 <<<grid, block>>> 启动语法
• 用 cudaMalloc / cudaMemcpy / cudaFree 三件套搬数据
• 明白为什么每个 CUDA 调用都必须检查返回码，错误检查宏怎么写

前置知识

第 1 章环境已搭好，能编译 C++ 程序。

核心概念

2.1 host vs device 心智模型

CUDA 程序里有两个执行世界：

graph LR
    H["host = CPU
跑 main()
分配 / 准备数据"] -- cudaMemcpy --> D["device = GPU
跑 kernel
巨量并行计算"]
    D -- cudaMemcpy --> H
    style H fill:#f3f1e8,stroke:#8b1538
    style D fill:#f3f1e8,stroke:#2f5d3a

函数靠修饰符区分：

关键代码

2.2 第一个 kernel: hello.cu

源码：hello.cu。

__global__ void hello_kernel() {
    printf("hello from thread (%d, %d) of block (%d, %d)\n",
           threadIdx.x, threadIdx.y,
           blockIdx.x,  blockIdx.y);
}

int main() {
    hello_kernel<<<2, 4>>>();   // 2 blocks × 4 threads = 8 lines
    KERNEL_CHECK();              // 等 kernel 跑完
}

启动语法 <<<...>>> 拆解

kernel<<<grid_dim, block_dim, shared_bytes, stream>>>(args);
//      ^^^^^^^^^  ^^^^^^^^^  ^^^^^^^^^^^^^  ^^^^^^
//      启动几个 block       每个 block 几个 thread
//                                     可选：动态 shared mem 大小
//                                                 可选：在哪个 stream 上

• grid_dim, block_dim：可以是 int（一维）或 dim3(x,y,z)（最多三维）。每维上限：block 内总线程 ≤ 1024。
• kernel 内可见的内置变量：threadIdx.{x,y,z} (本 block 内的线程号), blockIdx.{x,y,z} (本 grid 内的 block 号), blockDim, gridDim。

2.3 数据搬运: host_device_memcpy.cu

CUDA 上一切数据都活在显存里——你必须显式拷过去、再拷回来。

// 流程: 分配 → 拷入 → 计算 → 拷回 → 释放

float* d = nullptr;
CUDA_CHECK(cudaMalloc(&d, N * sizeof(float)));           // 1. 分配显存
CUDA_CHECK(cudaMemcpy(d, h, N*sizeof(float), cudaMemcpyHostToDevice));  // 2. H2D

scale_kernel<<<1, N>>>(d, 3.14f, N);                     // 3. 计算
KERNEL_CHECK();

CUDA_CHECK(cudaMemcpy(h, d, N*sizeof(float), cudaMemcpyDeviceToHost));  // 4. D2H
CUDA_CHECK(cudaFree(d));                                  // 5. 释放

本仓库提供了 RAII 封装 DeviceBuffer<T>（见 common/cuda_utils.h），少写一半代码：

DeviceBuffer<float> d(N);          // 析构时自动 free
d.copy_from_host(h);
scale_kernel<<<1, N>>>(d.ptr, 3.14f, N);
KERNEL_CHECK();
d.copy_to_host(h);

2.4 错误检查的三道关

CUDA 不像 C++ 抛异常——所有 API 都返回 cudaError_t，你不查就吞掉。错误源有三类：

所以仓库里的两个宏（common/cuda_utils.h）：

#define CUDA_CHECK(stmt)  do {                            \
    cudaError_t _e = (stmt);                              \
    if (_e != cudaSuccess) {                              \
        fprintf(stderr, "[CUDA] %s:%d: %s -> %s\n",       \
                __FILE__, __LINE__, #stmt,                \
                cudaGetErrorString(_e));                  \
        exit(1);                                          \
    }                                                     \
} while (0)

#define KERNEL_CHECK() do {                               \
    CUDA_CHECK(cudaPeekAtLastError());                    \
    CUDA_CHECK(cudaDeviceSynchronize());                  \
} while (0)

error_handling.cu 演示三种错误如何被这两个宏抓住，强烈建议自己跑一遍看输出。

运行结果

分别运行 hello、host_device_memcpy 与 error_handling，确认线程输出、CPU/GPU 对拍和预期错误模式与本章说明一致。

自检清单

练习题

1. 补全 exercises/01_axpy_starter.cu 中的 SAXPY kernel：y = a*x + y，N = 1M，与 CPU 版对拍。
2. 修改 error_handling.cu，在 bad_kernel 里写 *(int*)0 = 0（空指针解引用），观察 cudaDeviceSynchronize 的报错。
3. 把 scale_kernel 改成 fma_kernel：z[i] = x[i] * y[i] + bias。这是后面 MLP 的基本块。

2.7 工业实战：CUDA 调试工作流

教程里 printf + CUDA_CHECK 调试 100 行 kernel 够用。但当你 forward 跑了 30 个 kernel 然后某处输出 NaN，靠 printf 没用——必须上专业工具。

2.7.1 compute-sanitizer：CUDA 版 valgrind

compute-sanitizer（前身 cuda-memcheck）是定位 kernel 内 bug 的唯一可靠工具。四种工作模式：

# 1) memcheck — 越界、未初始化、非法指针
compute-sanitizer --tool memcheck ./my_app

# 2) racecheck — shared memory 数据竞争（漏 __syncthreads）
compute-sanitizer --tool racecheck ./my_app

# 3) initcheck — 用了未初始化的 device 内存
compute-sanitizer --tool initcheck ./my_app

# 4) synccheck — block 内 sync 不一致（部分 thread early-exit）
compute-sanitizer --tool synccheck ./my_app

典型输出：

========= Invalid __global__ write of size 4 bytes
=========     at 0x70 in matmul_kernel(float*, float*, float*, int, int, int)
=========     by thread (5,3,0) in block (1,0,0)
=========     Address 0x7f8e... is out of bounds
=========     Saved host backtrace up to driver entry point at error

编译时加 -lineinfo（不是 -g，会关优化）才能看到行号：

nvcc -O2 -lineinfo -arch=sm_80 mykernel.cu -o app
compute-sanitizer ./app

实践建议：发布前对每个 kernel 都跑一次 memcheck + racecheck，并在 CI 中保留代表性边界用例；检查耗时取决于 kernel 与输入规模。

2.7.2 cuda-gdb：源码级断点调试

支持 kernel 内打断点、单步、看变量。前提：编译用 -g -G（device 调试符号）；它会改变优化与时序，只用于调试，不能拿 debug build 做性能基线。

nvcc -g -G -arch=sm_80 mykernel.cu -o app_dbg
cuda-gdb ./app_dbg

(cuda-gdb) break mykernel.cu:42      # device 代码也能打断点
(cuda-gdb) run
(cuda-gdb) cuda thread (5,3,0) block (1,0,0)    # 切到特定 thread
(cuda-gdb) print x
(cuda-gdb) info cuda warps

VSCode + Nsight VSCode 插件能让你在 IDE 里点断点（非常推荐）。

2.7.3 常见数值 bug 诊断流程

症状：输出 NaN / Inf，定位顺序：

1. 在每个 kernel 后插临时 check_nan，二分定位是哪个 kernel 先产生 NaN
2. 看可疑 kernel 的输入：是输入就 NaN 了还是 kernel 算坏的？
3. kernel 算坏的话：检查除零（softmax 没减 max？rsqrt(0)？），log(负数)，exp(过大)

// 临时 NaN 探测 kernel — 加在 forward 各点对照
__global__ void nan_probe(const float* x, int n, int* found, const char* name) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n && (isnan(x[i]) || isinf(x[i])))
        atomicExch(found, 1);
}

// host: 跑完每个算子调一次
int h_found = 0, *d_found;
cudaMalloc(&d_found, 4);
cudaMemset(d_found, 0, 4);
nan_probe<<<...>>>(out, n, d_found, "after_layernorm");
cudaMemcpy(&h_found, d_found, 4, cudaMemcpyDeviceToHost);
if (h_found) fprintf(stderr, "NaN/Inf at after_layernorm!\n");

2.7.4 版本陷阱与诊断脚本

客户报"程序在新环境跑不起来"，9 成是版本不匹配。诊断脚本：

#!/usr/bin/env bash
# diag.sh — 在客户环境跑一次, 把输出粘给我
echo "=== 内核驱动 ==="
nvidia-smi --query-gpu=driver_version,cuda_version --format=csv,noheader

echo "=== CUDA Toolkit ==="
nvcc --version 2>/dev/null || echo "nvcc 不在 PATH"
ls -la /usr/local/cuda 2>/dev/null

echo "=== Python 侧 ==="
python -c "import torch; print('torch', torch.__version__, 'cuda', torch.version.cuda)" 2>/dev/null
python -c "import torch; print('available:', torch.cuda.is_available(), 'count:', torch.cuda.device_count())"

echo "=== libcudart ==="
ldconfig -p | grep libcudart

echo "=== 容器 ==="
[[ -f /.dockerenv ]] && echo "in docker" || echo "host"

2.7.5 生产 kernel 的 "不要" 清单

• 不要在生产 kernel 留 printf——device-side printf 有固定大小缓冲（默认 1 MB），高吞吐场景会丢日志，且严重拖慢吞吐
• 不要用 assert——device assert 失败直接 abort 整个 context，整张 GPU 都得重置
• 不要在 hot path 里 cudaDeviceSynchronize()——它会让所有 stream 都阻塞，破坏 overlap
• 不要裸用 cudaMalloc 在 hot path——每次 malloc 几百微秒。用 memory pool 或预分配
• 不要在 kernel 里 new / malloc（device-side heap，极慢且容量小）

2.8 研究前沿（2025-2026）：GPU kernel 开发的语言之战

"写 CUDA C++" 在 2026 已经不是唯一选择。下面是几个真正在生产中崛起的替代方案，每个都让你用更少的代码达到接近 CUDA 的性能。

2.8.1 Triton 3.0+（OpenAI，事实标准的"Python CUDA"）

Triton 用 Python-like 语法描述 tile-level 计算，编译器自动处理 block / warp / lane 分配。FlashAttention 早期版本、vLLM、SGLang 的大部分 kernel 都用 Triton 写。

import triton
import triton.language as tl

@triton.jit
def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):
    pid    = tl.program_id(axis=0)
    offsets = pid * BLOCK + tl.arange(0, BLOCK)
    mask    = offsets < n
    x = tl.load(x_ptr + offsets, mask=mask)
    y = tl.load(y_ptr + offsets, mask=mask)
    tl.store(out_ptr + offsets, x + y, mask=mask)

# Python 侧:
add_kernel[(triton.cdiv(n, 1024),)](x, y, out, n, BLOCK=1024)

2025-2026 关键更新：

• Triton 3.0+ 加 Hopper TMA / wgmma 完整支持
• Triton-Blackwell 支持 fp4 + TMEM
• Triton-CPU 让同一 kernel 跑 GPU 和 CPU
• TritonBench (Meta) — 100+ 个工业 kernel benchmark

2.8.2 ThunderKittens（Stanford, 2024-2025）— 100 行写 attention

ThunderKittens (TK) 是 Stanford Hazy Research 团队推出的 C++ 嵌入式 DSL，专为极简 Tensor Core 编程设计。一个完整的 FlashAttention v2 不到 100 行：

// ThunderKittens 风格 (简化伪代码)
using namespace kittens;
__global__ void attn_kernel(...) {
    auto Q_tile = make_tile<rt_fl, 16, 64>();   // tile abstraction
    auto K_tile = make_tile<rt_fl, 64, 16>();
    auto V_tile = make_tile<rt_fl, 16, 64>();

    load(Q_tile, Q_global);                      // 自动 cp.async + TMA
    for (int kt = 0; kt < n_kv_tiles; ++kt) {
        load(K_tile, K_global + kt);
        load(V_tile, V_global + kt);
        auto S = mm<trans_b>(Q_tile, K_tile);     // 自动 wgmma
        softmax_inplace(S);
        accumulate(O_tile, mm(S, V_tile));
    }
    store(O_global, O_tile);
}

ThunderKittens 提供较紧凑的 tile/warp abstractions，适合阅读和原型化；是否接近 CUTLASS 取决于硬件、shape、dtype 和实现，必须用同一 workload 测量。

2.8.3 CUTLASS 4.6.1 与 CuTe DSL（2026-07 快照）

CUTLASS 4.6.1 同时提供成熟的 CUDA C++ template abstractions 和 Python CuTe DSL。 CuTe 的 layout 代数可表达 tile / warp / instruction 映射，但不能据此声称所有 CUTLASS 路径都由 Python DSL 重写：

// CUDA C++ 中的 CuTe layout algebra；Python CuTe DSL 是另一条入口
using namespace cute;
auto thr_layout = make_layout(make_shape(_8{}, _32{}));   // 8x32 thread layout
auto data_layout = make_layout(make_shape(_128{}, _64{}),
                                make_stride(_64{}, _1{}));
auto tensor = make_tensor(ptr, data_layout);
auto thr_tile = local_partition(tensor, thr_layout, tid);
// 编译期就知道每 thread 拿哪些元素, 自动 swizzle / vectorize

学习曲线陡，但是 NVIDIA 长期主推方向。Blackwell 上的 TMEM GEMM 必须用 CuTe / CUTLASS 写。

2.8.4 Mojo / MAX（Modular，2024-2026）

Mojo 是 Python 超集 + 系统级性能，由 Chris Lattner（LLVM 创始人）团队设计。声称"Python 语法 × C++ 性能"，2025 年开始有 CUDA backend，能编译到 PTX。

定位介于 Triton 和 CUDA 之间，吸引不想学 C++ 但需要超过 Python 性能的研究者。生产部署目前还少，2026 年值得关注。

2.8.5 各方案对比

2026 务实选择建议：

• 学习 / 研究：先 CUDA C++ 打基础（本教程），然后转 Triton + ThunderKittens 提效率
• 新 kernel 原型：直接 Triton
• 极致性能：CUTLASS / CuTe（Blackwell 上几乎是必选）
• 跨平台：JAX/Pallas (TPU + GPU) 或者等 Mojo 成熟

2.8.6 AI 辅助 kernel 开发（2025 后兴起）

大模型开始能写出可用的 CUDA kernel：

• Sakana AI（2024 末）："AI CUDA Engineer" — LLM 自动生成 + benchmark CUDA kernel
• KernelBench（Stanford, 2025）— LLM 写 kernel 的标准评测
• Claude 3.5+/GPT-5 写 Triton 已经接近资深工程师水平（简单 kernel）

意味着2026 工程师的角色更偏 "kernel 设计 + 验证"，不再是"逐行写"。但读懂 CUDA、debug CUDA依然是不可替代的核心技能——这正是本教程的目标。

常见坑

• kernel 修改了数据但 host 读到旧值 → 忘记 cudaDeviceSynchronize（或 KERNEL_CHECK）。
• cudaErrorInvalidConfiguration → block 维度超 1024 或 shared mem 超限。
• cudaErrorIllegalAddress → kernel 内越界访问、空指针、misaligned 加载。用 cuda-memcheck 或 compute-sanitizer 定位行号。
• 整个程序"成功"但结果全 0 → 多半 kernel 没启动（参数错被无视）。永远加 KERNEL_CHECK。

2.10 CUDA 官方手册精讲（CUDA Programming Guide 13.2（核验：2026-07-20））

NVCC、Variable Specifiers、错误模型、Unified Memory

NVCC 编译流水线：一行命令背后的 5 步

当你跑 nvcc hello.cu -arch=sm_80 -o hello，nvcc 不是一个真编译器，而是一个编译器驱动（compiler driver）——它把 .cu 文件按 host 和 device 两条线分别处理，最后链回一个可执行：

    graph TD
        Src["hello.cu
(host + device 混编)"]
        Sep["nvcc 前端: 分离 host / device 代码"]
        HostC["host C++ 部分"]
        DevC["device C++ 部分 (__global__ / __device__)"]
        Hostcc["host 编译器
(gcc / clang / msvc)"]
        GpuCC["nvcc 自带 GPU 前端"]
        PTX["PTX (compute_80)"]
        Ptxas["ptxas 汇编器"]
        Cubin["cubin (sm_80)"]
        Fatbin["fatbin 容器"]
        Hostobj["host .o (含 fatbin 段)"]
        Linker["host linker"]
        Exe["可执行文件 hello"]

        Src --> Sep
        Sep --> HostC
        Sep --> DevC
        DevC --> GpuCC
        GpuCC --> PTX
        PTX --> Ptxas
        Ptxas --> Cubin
        PTX --> Fatbin
        Cubin --> Fatbin
        Fatbin --> Hostobj
        HostC --> Hostcc
        Hostcc --> Hostobj
        Hostobj --> Linker
        Linker --> Exe

        style Src fill:#f3f1e8,stroke:#8b1538
        style Exe fill:#f3f1e8,stroke:#2f5d3a
    

5 步详解

1. 分离：nvcc 前端扫描 .cu 文件，把 __global__ / __device__ 函数和它们引用的代码归为 device 代码，其余是 host 代码。
2. device 编译：device 代码经 nvcc 内置的 GPU C++ 前端编成 PTX（虚拟 ISA，文本格式）。
3. PTX → cubin：ptxas 把 PTX 汇编成对应 SM 的 cubin（二进制）。-arch=sm_80 决定这步的目标。
4. 打包 fatbin：把所有目标的 cubin + 一份 PTX 塞进 fatbin 容器，作为一个数据段嵌入 host 的 .o 文件里。
5. host 编译 + 链接：nvcc 把 host 代码交给系统 host 编译器（默认 gcc，Windows 上 MSVC），最后用系统 linker 链成可执行。可执行文件里既有 CPU 机器码，也有那段 fatbin。

看清楚每一步

# 1) 看完整的内部命令链
    nvcc -v hello.cu -arch=sm_80 -o hello 2>&1 | head -30

    # 2) 保留中间文件 (PTX, cubin, .o)
    nvcc -keep hello.cu -arch=sm_80 -o hello
    ls hello.*    # → hello.ptx hello.cubin hello.cudafe1.cpp ...

    # 3) 直接 dump 出可执行里的 PTX / cubin
    cuobjdump --dump-ptx hello
    cuobjdump --dump-sass hello   # SASS = 真实硬件指令 (cubin 反汇编)

常用的 nvcc 标志

nvcc 不是必需的

NVRTC 库（运行时编译）和 LLVM/Triton 等替代前端也能产出 PTX，最终都送进 Driver 的 JIT。本教程只用 nvcc，因为它最直接、错误信息最清晰。

Variable Specifiers：变量住在哪由你说了算

CUDA 给变量声明提供 4 个修饰符，告诉编译器这个静态变量该放到哪种内存里。本节先建立全图，后续章节会逐个深入用法。

各修饰符的示例

// 1) __constant__ — 全 thread 都读同一个值，速度快（自带 cache）
    __constant__ float attn_scale;   // 64 KB constant memory 内
    __global__ void attn_kernel(...) {
        // 所有 thread 同时读 attn_scale，硬件会广播
        float s = attn_scale;
        /* ... */
    }
    // host 侧写入:
    float h_scale = 1.0f / sqrtf(64.0f);
    cudaMemcpyToSymbol(attn_scale, &h_scale, sizeof(float));

    // 2) __device__ — 跨 kernel 持久的 global memory 变量
    __device__ int g_counter;        // 全局，所有 kernel 都看得见
    __global__ void inc_kernel() {
        atomicAdd(&g_counter, 1);
    }

    // 3) __managed__ — Unified Memory, host 和 device 都能直接访问
    __managed__ float g_loss;
    __global__ void loss_kernel(...) { g_loss = compute(); }
    int main() {
        loss_kernel<<<...>>>(...);
        cudaDeviceSynchronize();
        printf("loss = %f\n", g_loss);   // host 直接读, 不需 memcpy
    }

    // 4) __shared__ — block 内共享, 通常在 kernel 体内声明
    __global__ void reduce_kernel(float* x, float* out) {
        __shared__ float smem[256];
        smem[threadIdx.x] = x[blockIdx.x * 256 + threadIdx.x];
        __syncthreads();
        // ... 接 reduction
    }

函数修饰符的完整版

当前 2.1 节列了 4 个函数修饰符，再补一个常见的 inline 写法：

用 __CUDA_ARCH__ 分辨编译路径

写 __host__ __device__ 函数时，有时希望 GPU 上走 fast path、CPU 上走稳健 path。用预定义宏：

__host__ __device__
    inline float fast_exp(float x) {
    #ifdef __CUDA_ARCH__
        // device 编译路径: 用 GPU 内置快速近似指令
        return __expf(x);
    #else
        // host 编译路径: 标准 libm
        return ::expf(x);
    #endif
    }

    // 还可以按 CC 分:
    #if __CUDA_ARCH__ >= 900
        // 用 Hopper 才有的 wgmma
    #elif __CUDA_ARCH__ >= 800
        // 用 Ampere 的 cp.async
    #endif

错误模型深入：同步 vs 异步、sticky、CUDA_LOG_FILE

2.4 节列了三类错误源，这里把"错误状态机"讲透——能让你看懂为什么 production CUDA 代码满屏 CUDA_CHECK。

1) 每个 host 线程都有自己的 error state

CUDA Runtime 给每个 host 线程维护一个 cudaError_t 状态（这就是"per-thread error state"）。每次 API 调用：

• 返回本次操作的错误码（如果是 cudaSuccess 就是成功）
• 如果失败，还会写到该 host 线程的 error state 里
• 错误状态不会自己清除，必须用 cudaGetLastError() 主动读取并清零

2) Sticky error：一颗烂葡萄毁一筐

bad_kernel<<<...>>>(); // 这里越界, 但没检查 → error state = invalid memory access

    cudaMalloc(&p, 100);    // 这行返回 invalidMemoryAccess (其实不是 malloc 的错!)
    cudaMemcpy(...);         // 这行也返回 invalidMemoryAccess
    cudaFree(p);             // 还是返回 invalidMemoryAccess
    // 最后你以为 cudaFree 失败了, 实际是 bad_kernel 的锅

3) cudaGetLastError vs cudaPeekAtLastError

// 标准的 kernel 后双重检查
    my_kernel<<<grid, block>>>(...);
    CUDA_CHECK(cudaGetLastError());          // 抓 launch 配置错误 (block 超 1024 等)
    CUDA_CHECK(cudaDeviceSynchronize());     // 抓 kernel 运行时错误 (越界, 除零等)

4) 同步错误 vs 异步错误：什么时候才"暴露"

CUDA API 分两类，错误暴露时机不同：

5) CUDA_LOG_FILE：不写代码也能抓错误

CUDA 13.1+ 引入 CUDA_LOG_FILE 环境变量：所有错误自动写到文件，含完整原因描述。比应用代码自己的错误信息详细得多。

# 把错误日志写到 cudaLog.txt
    env CUDA_LOG_FILE=cudaLog.txt ./my_app

    # 直接打到 stderr (CI 友好)
    env CUDA_LOG_FILE=stderr ./my_app

    # 看日志
    $ cat cudaLog.txt
    [12:46:23.854][137216133754880][CUDA][E] One or more of block dimensions of
    (4096,1,1) exceeds corresponding maximum value of (1024,1024,64)
    [12:46:23.854][137216133754880][CUDA][E] Returning 1 (CUDA_ERROR_INVALID_VALUE)
    from cuLaunchKernel

对比应用层的 cudaGetErrorString 只输出 invalid argument 这种笼统话，CUDA_LOG_FILE 直接告诉你"是 4096 这个 block 维度超了 1024 上限"，定位精度跳一个档次。

6) cudaErrorNotReady 不算错

查询异步操作状态用 cudaStreamQuery / cudaEventQuery，如果操作还没完成，它返回 cudaErrorNotReady。这不是错误——只是"还在跑"的标志。cudaGetLastError 和 cudaPeekAtLastError 都会跳过它，不写入 error state。第 8 章 async 章节会大量用到。

7) 调试技巧：CUDA_LAUNCH_BLOCKING=1

异步 kernel 让错误来源很难定位（错误暴露时离真凶可能有几十个 API 调用）。设这个环境变量让每个 kernel 启动都立即同步，把"哪个 kernel 错的"暴露在错误那一行：

env CUDA_LAUNCH_BLOCKING=1 ./my_app

代价：会序列化原本可并行的 stream work。仅调试用，不应留在生产热路径。

CUDA Runtime 初始化：第一次 cudaMalloc 为什么慢

第一次跑 CUDA 程序的人常困惑：cudaMalloc(1024) 这么小的分配，竟然要 100-500 毫秒？因为它顺带做了整个 CUDA Runtime 的初始化。

"primary context" 是什么

CUDA Runtime 在每个 device 上维护一个 primary context（主上下文）。它第一次需要用到该 device 时才创建——通常是你的第一个 CUDA API 调用。创建 primary context 包括：

1. 把 device 代码（fatbin 里所有 cubin / PTX）JIT 编译并加载到 GPU 内存
2. 初始化 device 的 memory allocator
3. 建立 host-device 之间的通信通道
4. 初始化 stream 0（legacy default stream）

这一切对你完全透明，但首次加载/JIT 可能明显慢于 steady-state；开销随 cubin、driver 与缓存状态变化，同一进程通常只发生一次。

什么时候触发初始化？

性能测量时的坑

// 错: 把初始化时间算到 kernel 时间里
    GpuTimer t;
    t.start();
    my_kernel<<<...>>>(...);    // <-- 第一次调用 = runtime 初始化 + JIT 编译 + 真正 kernel
    cudaDeviceSynchronize();
    float ms = t.stop();          // 可能是 500 ms, 但 kernel 本身只用了 0.1 ms

    // 对: 先"预热"一次, 再计时
    my_kernel<<<...>>>(...);    // warmup: 触发初始化 + JIT
    cudaDeviceSynchronize();

    t.start();
    for (int i = 0; i < 100; ++i) my_kernel<<<...>>>(...);
    cudaDeviceSynchronize();
    float avg_ms = t.stop() / 100.0f;

本教程后续所有性能测量都遵守"warmup + N 次平均"的模式。第 8 章性能分析会反复强调。

显式初始化（CUDA 12.0+）

// 显式初始化指定 device 的 primary context
    CUDA_CHECK(cudaInitDevice(0, /*flags=*/0, /*deviceFlags=*/0));
    // 现在第一次 cudaMalloc 不再触发初始化, 时间可控

cudaDeviceReset 的双刃剑

cudaDeviceReset() 销毁当前 device 的 primary context，所有 device 内存、stream、event 都会被释放。之后再调用 CUDA API 会自动新建 context，相当于回到"第一次启动"状态。

显式内存 vs Unified Memory：何时选哪种

当前 2.3 用 cudaMalloc + cudaMemcpy + cudaFree 显式管理 device 内存。CUDA 还有另一种选择——Unified Memory（统一内存）：分配出来的指针 host 和 device 都能直接用，CUDA Driver 替你在背后搬运。

两种风格的代码对比

显式版本（本教程主流）

// 完整流程: 5 步
    float* h_x = (float*)malloc(N * sizeof(float));    // host buffer
    init(h_x, N);

    float* d_x;
    CUDA_CHECK(cudaMalloc(&d_x, N * sizeof(float)));   // 1. 显式 alloc on device
    CUDA_CHECK(cudaMemcpy(d_x, h_x, N*sizeof(float),
                          cudaMemcpyHostToDevice));    // 2. H2D copy

    kernel<<<grid, block>>>(d_x, N);                   // 3. kernel 用 d_x
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());

    CUDA_CHECK(cudaMemcpy(h_x, d_x, N*sizeof(float),
                          cudaMemcpyDeviceToHost));    // 4. D2H copy
    CUDA_CHECK(cudaFree(d_x));                          // 5. free
    free(h_x);

Unified Memory 版本

// 短得多
    float* x;
    CUDA_CHECK(cudaMallocManaged(&x, N * sizeof(float)));   // 一次分配
    init(x, N);                                              // host 直接初始化

    kernel<<<grid, block>>>(x, N);                          // device 直接用
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());                     // ← 必须 sync 后 host 才能安全读

    printf("x[0] = %f\n", x[0]);                             // host 直接读
    CUDA_CHECK(cudaFree(x));

取舍

Unified Memory 还有 4 种风格

从 PDF Sec 2.4.2 起，Unified Memory 在不同硬件上行为差异很大：

查当前 device 支持哪种：

int concurrent, pageable, hw_coherent;
    cudaDeviceGetAttribute(&concurrent, cudaDevAttrConcurrentManagedAccess, 0);
    cudaDeviceGetAttribute(&pageable,   cudaDevAttrPageableMemoryAccess, 0);
    cudaDeviceGetAttribute(&hw_coherent, cudaDevAttrPageableMemoryAccessUsesHostPageTables, 0);

    if (concurrent && pageable && hw_coherent)  puts("Full + hardware coherence (Grace)");
    else if (concurrent && pageable)             puts("Full + software coherence (HMM)");
    else if (concurrent)                          puts("Full but only cudaMallocManaged");
    else                                          puts("Limited (Windows/WSL/Tegra)");

Page-locked Host Memory：cudaMallocHost 的故事

本教程 1.4 节用了 cudaMallocHost，但没解释为啥。页锁定内存（pinned memory）有 3 个好处：

1. 异步拷贝必需：cudaMemcpyAsync 要求源/目标内存是 pinned 的，否则会退化为同步
2. 同步拷贝也更快：DMA 直接读，不经过 OS 内核 staging buffer
3. 能 map 到 device 地址：kernel 可以直接访问 host 内存（zero-copy）

4 个相关 API：

下一章导览

第 3 章我们把"启动 N 个线程并行处理 N 个元素"系统化：线程模型与索引。重点是 1D/2D 索引计算，以及处理"N 不是 block size 整数倍"的边界条件。