CUDA 编程入门：第一个 kernel，从 SAXPY 建立 host/device 心智模型

更新时间：2026-09-25。学会 AI 加速器为什么比 CPU 快之后，下一步自然是亲手给 GPU 写一段代码。GPU 不是一个「更快的 CPU」，它有自己独立的内存和成千上万个并发线程，直接把 C 语言的 for 循环搬过去跑不起来。本文是「集成电路大类 · 芯片设计」**入门层（beginner）**的一篇，用 CUDA 生态里最经典的例子 SAXPY 向量加法，带你走完从分配内存到拿回结果的完整闭环。读完你会理解 __global__、cudaMalloc、<<<grid, block>>> 这些怪语法各自在干什么；GPU 为什么走吞吐路线见 GPU 与 CPU 的取舍，线程层次的深入展开见 CUDA 线程层次。

本文要回答的问题

CUDA 程序为什么分 host 和 device 两边，代码里哪些在 CPU 跑、哪些在 GPU 跑？
__global__ 标记的 kernel 函数和普通 C 函数有什么不同，怎样被调用？
成千上万个线程执行同一份 kernel，每个线程怎么知道自己该处理数组的哪个元素？
数据一开始在 CPU 内存里，怎样送到显卡、算完又怎样拿回来？
用 nvcc 怎样编译，CUDA 出错时为什么常常「不报错」，应该怎样检查？

一、心智模型：host 与 device 是两台协作的机器

写 CUDA 之前先接受一个事实：你的程序同时运行在两台机器上。host（主机） 是 CPU 加它的系统内存（DRAM），负责启动程序、准备数据、调度 GPU；device（设备） 是 GPU 加它自己的显存（global memory），负责大规模并行计算。两边通过 PCIe（或 NVLink）总线相连，各自有独立的地址空间——CPU 的指针和 GPU 的指针互不通用，数据必须显式搬运。

对应到代码，一个 .cu 源文件里混着两类函数：不带修饰的普通 C/C++ 函数在 host 上由 CPU 执行；用 __global__ 修饰的函数叫 kernel（核函数），它在 host 端被调用、却在 device 上由成千上万个线程同时执行。kernel 必须返回 void，参数通常是指向显存的指针。

把一个熟悉的 CPU 串行程序和 CUDA 概念逐行对照，心智模型就立起来了：

CPU 串行代码里的概念	CUDA 里的对应物	一句话差别
`for (i=0; i<n; i++)` 顺序遍历	一个 grid 里铺开 n 个 thread	时间上串行变成空间上并行
循环变量 `i` 取 0,1,2…	`blockIdx.x*blockDim.x+threadIdx.x`	每个线程的编号是常量，只算一个元素
普通函数调用 `f(args)`	`f<<<grid, block>>>(args)` 启动 kernel	尖括号描述线程布局，且启动立即返回
`float *p; malloc(...)`	`float *d_p; cudaMalloc(...)`	显存独立分配，host 不能直接解引用
数组直接赋值传参	`cudaMemcpy(...)`	跨总线按字节整块拷贝，要标方向
函数返回 = 结果已算完	kernel 启动返回 ≠ 执行完	kernel 是异步的，要靠同步点确认
程序崩溃有明确报错	kernel 非法访问常静默失败	每个 CUDA 调用都要检查错误码

补充一点背景：CUDA（Compute Unified Device Architecture，统一计算架构）是 NVIDIA 在 2007 年推出的并行计算平台和编程模型，所以这套代码只能在 NVIDIA 的 GPU 上跑；其他厂商的显卡有各自对应的生态（ROCm、Metal 等），概念相通但 API 不同。这种「CPU 加协处理器、各自有内存」的写法统称异构计算（heterogeneous computing）。为什么非要搞得这么分裂，不让 GPU 直接访问主机内存？物理上它们就是两颗芯片、两种内存颗粒，中间隔着总线；显存为高带宽做了专门设计，而系统内存要兼顾通用性和容量，二者延迟和带宽差一个数量级。显式搬运虽然啰嗦，却让程序员对「什么时候付了过桥费」心中有数。

记住这张表，后面所有新 API 都是在补「两台机器之间协作」的细节。本文反复出现的以 d_ 开头的指针（d_x、d_y）只是命名约定——提醒你这是 device 上的地址，编译器并不强制；但几乎所有 CUDA 代码库都遵守它，建议从第一天起照做。

二、任务：SAXPY，每个元素互相独立

我们要实现的是 BLAS 数学库中的经典算子 SAXPY（Single-precision A times X Plus Y）：给定标量 a 和两个长度为 n 的向量 x、y，计算

bash

y[i] = a * x[i] + y[i]    （i = 0, 1, …, n-1）

它的 CPU 版本只有一个循环：

// 在 host 上运行的普通 C 函数
void saxpy_cpu(int n, float a, const float *x, float *y) {
    for (int i = 0; i < n; ++i)
        y[i] = a * x[i] + y[i];
}

选它做第一个 CUDA 程序有两个理由。其一，计算只有一次乘法一次加法，注意力可以全部放在流程而不是算法上。其二，也是关键：第 i 个元素的结果只依赖 x[i]、y[i]，与其他下标毫无关系。这意味着 n 个元素可以毫无顾忌地同时计算，不需要线程之间通信或等待——这正是 GPU 最喜欢的「数据并行」形态。反过来，如果每一步都依赖上一步的结果（比如递推求和），就不能这样直接并行化。

例子里取 n = 1 << 20，即 1048576 个元素，约占 4 MB 浮点内存。为什么不用小一点的数组演示？因为 GPU 的优势要在数据量足够大时才显现：数据太少，启动 kernel 和跨总线拷贝的固定开销会盖过计算收益，结果可能比 CPU 还慢。这也提示了一条朴素的使用判据——先问搬运划不划算，再问要不要上 GPU。

三、写 kernel：每个线程认领一个下标

CUDA 版 SAXPY 不长，但每一行都有讲究：

cpp

// 在 device 上运行、由 host 启动的 kernel
__global__ void saxpy_kernel(int n, float a,
                             const float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // 全局线程编号
    if (i < n)                                      // 越界保护
        y[i] = a * x[i] + y[i];
}

__global__ 告诉编译器：这个函数编译成 GPU 指令，由 host 端通过 <<<...>>> 语法启动。一次启动会产生一大批线程，每个线程都执行这同一份代码，但各自看到的内置变量不同。CUDA 给每个线程准备了几个只读的内置变量：

threadIdx.x：我在当前 block 内排第几（0 到 block 大小减 1）；
blockIdx.x：我所在的 block 在 grid 内排第几；
blockDim.x：每个 block 有多少个线程；
此外还有 gridDim.x：本次启动一共有多少个 block，kernel 里偶尔用得到。

这些变量后面的 .x 是因为它们都是三维向量类型（uint3），还带 .y、.z 分量。处理一维数组时只用 .x；处理图像或体数据时，block 和 grid 可以排成二维、三维，线程直接拿到「行列坐标」，不用自己把一维编号再除来除去。本文先把一维情形吃透。

把 block 想象成固定载客量的班车：blockDim.x 是每车人数，blockIdx.x 是第几辆车，threadIdx.x 是车上第几号座位。要得到「全 grid 的连续座位号」，公式就是

bash

i = blockIdx.x * blockDim.x + threadIdx.x

例如每车 256 人，第 0 辆车的线程拿到 0..255，第 1 辆车拿到 256..511，严丝合缝。这个一维索引公式是 CUDA 里出现频率最高的一行，务必背下来。

if (i < n) 不是可有可无的。block 大小固定（常用 256），而 n 未必是 256 的整数倍。host 端算 block 数量时只能向上取整，于是最后一个 block 里总会多出几个线程，它们的 i 会越过数组末尾。不做判断就写 y[i]，结果就是非法显存访问，轻则结果错乱、重则程序崩溃。256 这个数字也不是语法规定：它是 32（warp 大小）的整数倍，又能让一个 block 占满较多 SM 资源，属于久经考验的默认值，深入讨论见 CUDA 线程层次。

四、main 函数：七步骨架与执行流程

host 端的 main 负责把整件事串起来，固定是七步：分配主机内存、分配显存、把数据拷过去、启动 kernel、把结果拷回来、释放所有内存。整体时序如下图。

图中步骤 4 和步骤 5 之间的时序关系是新手最大的认知坎：kernel<<<...>>> 这句的语义更像「把任务单扔给 GPU」，host 线程发出启动命令后立刻继续执行下一行，GPU 在后台并行地跑。步骤 6 的 cudaMemcpy 方向是 D2H，它在搬运前会隐式等待 kernel 完成，所以这个最简例子里结果不会错——但一旦你想在 host 端计时、或者启动后立刻做别的事，就必须理解这条异步规则。

下面是可编译运行的完整 main：

cpp

#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>

int main(void) {
    int n = 1 << 20;                 // 100 万个元素
    size_t bytes = n * sizeof(float);

    // 1. host 内存：准备输入
    float *x = (float*)malloc(bytes);
    float *y = (float*)malloc(bytes);
    for (int i = 0; i < n; ++i) { x[i] = 1.0f; y[i] = 2.0f; }

    // 2. device 内存：在显存开两块等大的空间
    float *d_x = NULL, *d_y = NULL;
    cudaMalloc(&d_x, bytes);
    cudaMalloc(&d_y, bytes);

    // 3. H2D：host → device
    cudaMemcpy(d_x, x, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, y, bytes, cudaMemcpyHostToDevice);

    // 4. 启动 kernel：总线程数向上取整保证覆盖 n
    int block = 256;
    int grid  = (n + block - 1) / block;
    saxpy_kernel<<<grid, block>>>(n, 2.0f, d_x, d_y);

    // 6. D2H：device → host（此处会等步骤 5 完成）
    cudaMemcpy(y, d_y, bytes, cudaMemcpyDeviceToHost);

    // 7. 清理
    cudaFree(d_x);
    cudaFree(d_y);
    free(x);
    free(y);
    return 0;
}

几个细节值得停下来看。cudaMalloc 的用法和 malloc 不同：它接收指针的地址（二级指针），由运行时回填显存地址，所以要传 &d_x；它的第二个参数是字节数，不是元素个数。grid 大小用整数公式 (n + block - 1) / block 做向上取整，保证 grid × block ≥ n。kernel 参数里传的是 d_x、d_y 这类显存指针，而不是 host 的 x、y——传反了，kernel 里就是非法地址。

你可能会担心：100 万个线程，真的在同一瞬间全部执行吗？不需要。GPU 只把当时装得下的 block 放进 SM，跑完一批再补一批，这叫「随退随补」（LBM）式调度。正因为如此，kernel 代码不应该假设 block 的执行先后：block 0 不一定比 block 99 先跑。这个设计还有一个巨大好处——同一份代码，在高端卡上铺开几千个 block、在入门核显上少铺几批，都能正确运行，性能随硬件自动伸缩。这就是为什么 grid 大小按数据规模 n 倒推，而不是查询「这张卡有多少核」再硬编码。

五、数据怎样过桥：内存搬运全景

初学者常把 cudaMemcpy 理解成「赋值」，它其实更像一次跨设备的快递：参数依次是目的地、源地、字节数、方向枚举。整个搬运格局如下。

读这张图记住三件事。第一，kernel 只能访问显存，所以送数（H2D）和取结果（D2H）一次都不能省，这也是 PCIe 带宽常常成为小算子瓶颈的原因——数据量太小时，搬运开销可能盖过计算收益。给一个量级感受：PCIe 4.0 x16 的单向带宽约 32 GB/s，而中端显卡的显存带宽往往超过 500 GB/s，差十几倍，数据过桥确实是「慢车道」。第二，长度一律给字节数，n * sizeof(float)，给成元素个数是排名第一的高频笔误。第三，方向枚举有四种（H2D、D2H、D2D、H2H），源和目的与方向写得不一致，运行时会直接报错——前提是你检查了返回值。

顺便纠正一个常见的想当然：不要试图靠「让 GPU 直接看见主机数组」来省掉拷贝。普通 malloc 出来的是分页内存，GPU 无法稳定访问。进阶做法是用 pinned memory（锁页内存），即 cudaMallocHost 分配的、不会被操作系统换页的主机内存，它能让跨总线拷贝更快、还支持异步传输，但这属于后续优化话题；入门阶段老老实实走标准流程即可。

六、用 nvcc 编译、运行，以及错误检查

CUDA 源文件以 .cu 为后缀，编译器是随 CUDA Toolkit 安装的 nvcc。它是一个编译器驱动：把 host 代码交给系统的 C/C++ 编译器，把 device 代码编译成 GPU 的机器码（PTX 再到具体架构的 cubin）。最常用的编译命令只有一行：

bash

nvcc -O2 saxpy.cu -o saxpy    # 编译
./saxpy                       # Linux 下运行；Windows 上是 saxpy.exe

-O2 是常规优化等级。第一次使用前先跑 nvcc --version 确认工具链装好、再用 nvidia-smi 确认机器上确实有 NVIDIA GPU 且驱动正常。顺带一提，nvcc 默认会生成兼容若干代架构的代码；如果想为手头的显卡生成最优机器码，可以加 -arch=sm_80 这样的参数（数字对应架构代号），初学阶段无需深究，用默认值即可。

CUDA 排错最折磨人的一点是：kernel 内部出错往往不会让程序立刻崩溃。非法访问、启动参数非法，错误被记录在一个「粘滞」的错误状态里，如果没人查询，程序可能带着错误一路跑到结束。所以工程上有两个标配动作。

第一，用宏包装每一个 Runtime API 调用，检查返回的 cudaError_t：

cpp

#define CUDA_CHECK(call)                                                    \
  do {                                                                      \
    cudaError_t err__ = (call);                                             \
    if (err__ != cudaSuccess) {                                             \
      fprintf(stderr, "CUDA 错误 %s:%d: %s\n", __FILE__, __LINE__,         \
              cudaGetErrorString(err__));                                   \
      exit(EXIT_FAILURE);                                                   \
    }                                                                       \
  } while (0)

于是 cudaMalloc、cudaMemcpy 都写成 CUDA_CHECK(cudaMalloc(...))，出错时能看到文件名、行号和人类可读的原因。第二，kernel 启动本身没有返回值，启动后要紧接两行：

cpp

saxpy_kernel<<<grid, block>>>(n, 2.0f, d_x, d_y);
CUDA_CHECK(cudaGetLastError());      // 捕获启动配置错误（grid/block 非法等）
CUDA_CHECK(cudaDeviceSynchronize()); // 等 kernel 真正跑完，并捕获执行期错误

cudaGetLastError 专门接住启动那一刻的错误；cudaDeviceSynchronize 让 host 阻塞到 GPU 收工，执行期间的越界访问会在这里暴露。调试阶段强烈建议每个 kernel 后都加这两行，程序正确后再按需去掉同步以恢复异步性能。

最后，结果对不对要靠验证：初始化时让 x[i]=1.0f、y[i]=2.0f、a=2.0f，那么每个结果都应当是 4.0f。拷回后抽样（比如检查前若干项和最后一项）与 CPU 版本逐元素比较，误差在浮点容限内才算通过。只打印第一个元素是很多教程的偷懒做法，恰好漏掉最后一个 block 的越界问题。

七、初学者高频坑清单

以为 kernel 启动会等它跑完。<<<...>>> 是异步的，后面紧跟的 host 代码可能先执行。计时、连续启动、host/device 混算时要用 cudaDeviceSynchronize 显式划定等待点。
大小写成元素个数。cudaMalloc 和 cudaMemcpy 都要字节数 n*sizeof(float)。少乘 sizeof 是最常见的笔误，现象是部分数组内容莫名丢失。
忘记越界判断。grid 向上取整必然产生多余线程，kernel 里不写 if (i < n) 就是越界写。
host/device 指针混用。把 x 传进 kernel、或在 host 上读写 d_x，都是访问对面地址空间，必错。
拷贝方向写反。目的地址、源地址、cudaMemcpy... 枚举三者要一致；送数是 H2D，取结果是 D2H。
不检查错误。kernel 里的错误默认沉默，不调用 cudaGetLastError/同步检查，你可能在调试一个根本没成功执行的 kernel。
block 大小随手乱写。用不是 32 倍数的数字（比如 100、300）不会报错，但硬件仍按 32 线程的 warp 调度，会白白浪费槽位；初学一律用 256。

这一篇在知识体系里的位置

想理解 GPU 为什么需要这种海量线程的写法 → GPU 与 CPU 的吞吐取舍、AI 加速器入门；
想把本文的 grid/block/thread 索引学透（二维三维、warp、shared memory、占用率）→ CUDA 线程层次；
想看线程在硬件上具体怎样被调度执行 → GPU 架构与 warp 调度器。

总结

一句话总结：CUDA 编程的第一道门槛是建立 host/device 心智模型——程序同时跑在 CPU（host，配系统内存）和 GPU（device，配独立显存）这两台通过 PCIe 相连、地址空间互不相通的机器上，源码里普通 C 函数在 CPU 执行，__global__ 修饰、返回 void 的 kernel 在 GPU 上由成千上万个线程同时执行；经典入门算子 SAXPY（y[i]=a*x[i]+y[i]）之所以适合并行，是因为每个元素只依赖相同下标的输入、彼此完全独立，串行的 for 循环被替换成空间上铺开的 grid，每个线程用 CUDA 最核心的索引公式 i=blockIdx.x*blockDim.x+threadIdx.x 从 block 编号、block 大小、线程编号三个内置变量算出自己唯一的全局下标，并必须用 if(i<n) 挡住 grid 向上取整产生的越界线程；host 端 main 固定七步——malloc 准备主机数组、cudaMalloc（传指针地址、给字节数）开显存、cudaMemcpy 按 H2D 方向拷入、<<<grid,block>>> 启动（block 常用 256，grid 按 (n+block-1)/block 向上取整）、D2H 拷回、最后 cudaFree/free 清理；关键时序细节是 kernel 启动异步、立即返回控制权，而 D2H 拷贝会隐式等待 GPU 完成；工程上要用 CUDA_CHECK 宏检查每个 Runtime API 的返回码，并在 kernel 后调用 cudaGetLastError 与 cudaDeviceSynchronize 捕获否则会沉默的启动与执行期错误，再用 nvcc 一行命令 nvcc -O2 saxpy.cu -o saxpy 编译、抽样比对 CPU 结果验证；初学者最常见的坑就是忘记异步性、把字节数写成元素个数、漏掉越界判断、混用 host 与 device 指针、写反拷贝方向以及不查错误码——避开它们，第一个 CUDA 程序就真正跑通了。

入门

进阶

专家

CUDA 编程入门：第一个 kernel，从 SAXPY 建立 host/device 心智模型 ​

本文要回答的问题 ​

一、心智模型：host 与 device 是两台协作的机器 ​

二、任务：SAXPY，每个元素互相独立 ​

三、写 kernel：每个线程认领一个下标 ​

四、main 函数：七步骨架与执行流程 ​

五、数据怎样过桥：内存搬运全景 ​

六、用 nvcc 编译、运行，以及错误检查 ​

七、初学者高频坑清单 ​

这一篇在知识体系里的位置 ​

总结 ​

CUDA 编程入门：第一个 kernel，从 SAXPY 建立 host/device 心智模型

本文要回答的问题

一、心智模型：host 与 device 是两台协作的机器

二、任务：SAXPY，每个元素互相独立

三、写 kernel：每个线程认领一个下标

四、main 函数：七步骨架与执行流程

五、数据怎样过桥：内存搬运全景

六、用 nvcc 编译、运行，以及错误检查

七、初学者高频坑清单

这一篇在知识体系里的位置

总结