CUDA系统拆解-02-第一个CUDA程序：最小闭环与代码执行路径

本文是「CUDA系统拆解」系列第 02 篇。
系列导读：CUDA系统拆解-00-导读：从编程模型到 AI 推理系统的学习路线
上一篇：CUDA系统拆解-01-CUDA为什么存在：从计算负载到编程模型
下一篇：CUDA系统拆解-03-线程组织模型：grid、block、thread 到底在表达什么

1. 这篇解决什么问题

• CUDA 程序为什么天然分成 host 和 device 两部分。
• kernel 和普通函数到底有什么区别。
• <<<grid, block>>> 这套 launch 语法在表达什么。
• cudaMalloc、cudaMemcpy、cudaDeviceSynchronize、cudaFree 分别在做什么。
• 一个最小 CUDA 程序从 CPU 发起，到 GPU 执行，再把结果取回，完整链路到底怎么走。

2. 先记住的核心结论

• CUDA 程序不是“整个程序都跑在 GPU 上”，而是 CPU 控制、GPU 计算。
• host 负责准备数据、申请显存、提交任务、同步和取回结果；device 负责并行执行 kernel。
• kernel 不是普通函数，它是 GPU 上的大规模并行执行入口。
• <<<grid, block>>> 不是语法装饰，而是在定义线程组织方式。
• 最小闭环就是：准备数据 -> cudaMalloc -> cudaMemcpy(H2D) -> launch kernel -> 同步 -> cudaMemcpy(D2H) -> 释放资源。
• 理解这条最小执行链，后面学 warp、内存层级、stream、profiling、推理引擎才不会断层。

3. 正文讲解

3.1 CUDA 程序不是“跑在 GPU 上”，而是“CPU 控制 GPU 干活”

初学 CUDA 最容易犯的错，就是把 .cu 文件理解成“整段程序都在 GPU 上运行”。

更准确的理解是：

• 主控逻辑仍然在 CPU 上
• GPU 是被 CPU 调用的加速器
• CUDA 程序本质上是 host code + device code

其中：

• host code 运行在 CPU 上，负责准备数据、申请 GPU 显存、数据拷贝、启动 kernel、同步和结果校验
• device code 运行在 GPU 上，主要就是 kernel

你应该先建立下面这张脑图：

CPU(host)
  ├─ 准备数据
  ├─ 申请 GPU 内存
  ├─ 把输入拷到 GPU
  ├─ 启动 kernel
  ├─ 等待或同步
  └─ 把结果取回

GPU(device)
  └─ 并行执行 kernel

这就是 CUDA 最小闭环的骨架。

3.2 为什么必须分 host 和 device

这不是 CUDA 故意把编程搞复杂，而是硬件现实决定的。

CPU 和 GPU 是两类不同处理器：

• 执行模型不同
• 擅长的问题不同
• 内存空间通常不同
• 调度方式不同

CPU 更适合：

• 复杂控制流
• 分支和逻辑调度
• 系统调用、文件、网络、IO

GPU 更适合：

• 大规模规则并行
• 高吞吐张量计算
• 很多线程做相似操作

所以 CUDA 的设计很明确：

CPU 负责控制和调度，GPU 负责高吞吐并行计算。

这也是为什么推理框架看起来在 CPU 侧运行，但底层真正重计算的地方会落到 CUDA kernel 或 CUDA 库。

3.3 kernel 是什么，和普通函数有什么区别

普通 C++ 函数的调用语义很简单：CPU 调，CPU 执行。

CUDA kernel 不是这样。

例如：

__global__ void vecAdd(const float* A, const float* B, float* C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

这里的 __global__ 表示：

• 这是一个 kernel
• 从 host 侧发起
• 在 device 上执行

它和普通函数最大的区别，不在“写法长得不一样”，而在执行语义不一样：

• 普通函数通常执行一次
• kernel 会被 GPU 扩展成大量线程并行执行

所以你写 kernel 时，写的不是“整个任务如何串行做完”，而是：

单个线程该做什么

这就是 CUDA 思维和普通 C++ 思维的第一个分界线。

3.4 最小代码骨架长什么样

下面是一个足够小、但已经完整体现执行链的例子：

#include <cuda_runtime.h>

__global__ void vecAdd(const float* A, const float* B, float* C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    int N = 1024;
    size_t bytes = N * sizeof(float);

    float *h_A, *h_B, *h_C;   // host memory
    float *d_A, *d_B, *d_C;   // device memory

    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);

    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

    cudaDeviceSynchronize();
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
}

这段代码虽然简单，但已经包含了最核心的动作：

1. 分配 device memory
2. 把输入从 host 拷到 device
3. 用 <<<grid, block>>> 启动 kernel
4. 等待 GPU 完成
5. 把结果从 device 拷回 host
6. 释放 device memory

你后面看到的大多数 CUDA 程序，骨架都还是这套流程，只是更复杂、更并发、或者做了更多优化。

3.5 <<<grid, block>>> 到底在表达什么

这一行是 CUDA 最有辨识度的语法：

vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

它表示：

用某种线程组织方式去启动这个 kernel

其中：

• blockSize：每个 block 有多少线程
• gridSize：总共有多少个 block

所以总线程数大致是：

gridSize * blockSize

为什么 CUDA 不直接让你说“启动 N 个线程”，而是要分 grid 和 block？

因为 block 不是单纯分组，它还是：

• 调度单位
• 协作单位
• 后面 shared memory 和同步的基本作用范围

所以 grid / block 是执行模型的一部分，不是装饰语法。

3.6 为什么要先算 idx，还要写边界判断

这两行在 CUDA 代码里极其常见：

int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) { ... }

第一行的作用是算全局索引，也就是“当前线程该处理哪一份数据”。

这里用到的几个量分别表示：

• threadIdx.x：线程在 block 内的局部编号
• blockIdx.x：block 在 grid 内的编号
• blockDim.x：每个 block 的线程数

第二行边界判断也几乎总是需要。

原因很简单：线程组织通常按 block 对齐，而数据规模未必正好整除 block size。

例如：

• N = 1000
• blockSize = 256
• gridSize = (1000 + 255) / 256 = 4
• 总线程数 = 4 * 256 = 1024

这时会多出 24 个线程。如果不写 if (idx < N)，这些线程就会访问越界内存。

所以这套写法不是模板习惯，而是 CUDA 执行模型的自然结果：

• 先按硬件友好的方式组织线程
• 再用边界判断裁掉多出来的尾巴

3.7 cudaMalloc、cudaMemcpy、cudaDeviceSynchronize 分别在做什么

这几个 API 是最小闭环里的关键节点。

cudaMalloc

• 在 GPU 的 global memory 上申请空间
• 对应的是 device memory

这里必须明确：

• h_A / h_B / h_C 在 host memory
• d_A / d_B / d_C 在 device memory

这两边默认不是同一块空间。

cudaMemcpy

• 负责在 host 和 device 之间搬数据
• 常见方向是 HostToDevice 和 DeviceToHost

这一步非常重要，因为在很多真实系统里，搬数据本身就是性能成本。

cudaDeviceSynchronize

• 让 CPU 等待，直到 GPU 上之前提交的任务完成

这也说明一个关键事实：

CPU 提交 kernel，不等于 CPU 会自动等 GPU 算完

理解这点，后面才能继续理解 stream、异步执行和 overlap。

3.8 从 .cu 到实际执行，到底发生了什么

这一层先建立高层直觉，不展开到编译细节。

一个最小 CUDA 程序的大致路径是：

1. 你写 .cu 文件，里面同时包含 host 代码和 device 代码。
2. 编译阶段会把 host 部分处理成 CPU 可执行部分，把 device 部分处理成给 GPU 使用的代码。
3. 运行时由 CPU 执行主程序。
4. CPU 通过 Runtime API 申请 device memory、搬数据、提交 kernel。
5. GPU 接到 launch 请求后，按 grid / block / thread 组织并行执行。
6. 结果写回 device memory。
7. CPU 同步后，再把结果取回 host。

这就是“从 .cu 到实际执行”的最小直觉版路径。后面讲 PTX / SASS / 编译链路时，会把这一层再展开。

3.9 用一条执行时间线把这篇串起来

把上面的内容连成时间线，会更清楚：

CPU(host)
  1. 准备输入数据
  2. cudaMalloc 申请 d_A, d_B, d_C
  3. cudaMemcpy: h_A -> d_A
  4. cudaMemcpy: h_B -> d_B
  5. launch vecAdd<<<grid, block>>>
  6. cudaDeviceSynchronize
  7. cudaMemcpy: d_C -> h_C
  8. 检查结果并释放资源

GPU(device)
  1. 接收输入数据
  2. 接收 kernel launch 请求
  3. 创建 grid / block / thread
  4. 每个线程计算自己的 idx
  5. 满足 idx < N 的线程执行计算
  6. 结果写回 d_C

这条时间线很重要。因为后面所有复杂 CUDA 程序，本质上都是把这条链：

• 拉长
• 并发
• 重叠
• 做更精细的资源管理

但骨架不会变。

4. 和 AI 推理的关系

这篇虽然很基础，但它和 AI 推理的关系很直接。

推理框架表面上可能在做：

• 请求处理
• batching
• cache 管理
• 调度

但底层仍然要重复执行同一套最小闭环：

• 准备输入 tensor
• 管理显存
• 启动 CUDA kernel 或 CUDA 库
• 同步或继续推进下一步

所以后面你会看到很多推理优化，其实都在压这条链上的某个成本：

• 减少 cudaMemcpy
• 减少 kernel 次数
• 降低 launch overhead
• 让多个阶段重叠执行
• 让数据尽量长时间留在 GPU 上

也就是说，这篇讲的不是“教学玩具”，而是所有推理执行链的最小骨架。

5. 常见误区

• kernel 就是普通函数。不对，普通函数通常是 CPU 调 CPU 执行；kernel 是 host 发起、device 并行执行。
• 写了 CUDA 代码，数据就自动在 GPU 上。不对，host memory 和 device memory 通常需要显式管理和拷贝。
• <<<grid, block>>> 只是语法装饰。不对，它决定了线程组织方式，也是后续性能分析的入口。
• launch 之后 CPU 一定会等 GPU 算完。不对，提交和完成通常不是同一件事。
• 会写一个向量加法就等于懂 CUDA。不对，真正重要的是能解释这条执行链为什么这样设计、代价在哪、以后怎么优化。

6. 复习自测

• 为什么 CUDA 程序天然要分成 host 和 device？
• kernel 和普通函数最本质的区别是什么？
• <<<grid, block>>> 到底在表达什么？
• 为什么几乎总要写 idx 计算和 if (idx < N)？
• cudaMalloc、cudaMemcpy、cudaDeviceSynchronize 分别在解决什么问题？
• 一个最小 CUDA 程序从 CPU 发起到 GPU 执行再返回结果，完整路径是什么？
• 为什么说推理系统的很多优化，本质上都是在压这条最小执行链上的某个成本？

---

系列导航

• 导读：CUDA系统拆解-00-导读：从编程模型到 AI 推理系统的学习路线
• 上一篇：CUDA系统拆解-01-CUDA为什么存在：从计算负载到编程模型
• 下一篇：CUDA系统拆解-03-线程组织模型：grid、block、thread 到底在表达什么
• 延伸阅读：GPU系统拆解-05-CUDA 执行映射：抽象模型如何落到真实硬件
• 延伸阅读：CUDA实践-01-基础概念与分块