第 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 演示三种错误如何被这两个宏抓住，强烈建议自己跑一遍看输出。

2.5 自检清单

2.6 练习

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，能抓出 80% 的隐藏 bug。CI 流水线里也会跑（虽然慢 10-100×）。

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

支持 kernel 内打断点、单步、看变量。前提：编译用 -g -G（device 调试符号），会关闭所有优化，性能掉 10×+，仅用于调试。

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);
}

性能逼近手写 CUTLASS（90%+ peak）。是 2025 后教研用的事实标准。FA v3 重写、各种新 attention 变体都用 TK 原型化。

2.8.3 CUTLASS 3.x + CuTe DSL（NVIDIA 官方）

CUTLASS 3.5+ 全面用 CuTe（Layout-of-Tensor）DSL 重写。CuTe 把"tile / warp / instruction"统一成 layout 代数，能在编译期推导出所有索引和访问模式：

// CuTe 风格 (Python 绑定也快出了, 2025)
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依然是不可替代的核心技能——这正是本教程的目标。

2.9 常见坑

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

下一步

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