<<<CUDA C++ 学习路线grid · block · warp · lane
EN中文
阶段 1 · 起步01 / 15约 18 分钟cuda/01_hello.cu

第一个 kernel：让 CPU 指挥 GPU 干活

host 与 device 的分工、__global__ 的含义、kernel launch 为什么是异步的

学完这一课你会
  • 分清 host（CPU）代码和 device（GPU）代码，知道各自跑在哪
  • 看懂 __global__、__device__、__host__ 三个函数限定符的区别
  • 理解 kernel launch 是异步的，以及为什么必须同步
  • 能用 nvcc 编译并跑通第一个 CUDA 程序

写 CUDA 之前，先把一件事想清楚：你写的是一份源码，但它会被编译成两套二进制。一套跑在 CPU 上（host），负责分配内存、搬数据、发命令；另一套跑在 GPU 上（device），负责真正的并行计算。nvcc 的核心工作就是把这两部分从同一个 .cu 文件里拆出来，分别交给主机编译器和 GPU 后端。

   Host (CPU)                             Device (GPU)
   ──────────                             ────────────
   cudaMalloc      ──── 分配显存 ────▶     [ global memory ]
   cudaMemcpy H2D  ──── 拷贝输入 ────▶     [ a, b ]
   kernel<<<G,B>>> ──── 提交任务 ────▶     [ SM0 SM1 SM2 ... ]
        │                                        │
   （立即返回，CPU 继续跑）                    （并行执行 kernel）
        │                                        │
   cudaDeviceSynchronize  ◀──── 等待完成 ────────┘
   cudaMemcpy D2H  ◀──── 取回结果 ────       [ c ]
   cudaFree        ──── 释放显存 ────▶
一次 kernel 调用的完整生命周期

三个函数限定符

限定符在哪里执行谁可以调用典型用途
__global__GPUhost（也可被 device 动态并行调用）kernel 入口，必须返回 void
__device__GPU只能被 GPU 代码调用kernel 内部的辅助函数
__host__CPUCPU普通函数，默认就是它
__host__ __device__两边都编译一份两边都能调共用的小工具函数，如 clamp
01_hello.cu — 每个线程报出自己的身份
1#include <cstdio>2 3// __global__ = 在 GPU 上执行、由 CPU 发起调用的函数。4// kernel 必须返回 void，结果只能通过显存带出来。5__global__ void hello() {6    int tid = blockIdx.x * blockDim.x + threadIdx.x;7    printf("block %d / thread %d -> global id %d\n",8           blockIdx.x, threadIdx.x, tid);9}10 11int main() {12    hello<<<2, 4>>>();   // 2 个 block，每个 block 4 个线程，共 8 个线程13 14    // kernel launch 是异步的：这一行之前 CPU 根本没等 GPU。15    // 少了这次同步，main 直接返回，进程退出，你可能一行输出都看不到。16    cudaDeviceSynchronize();17 18    cudaError_t err = cudaGetLastError();19    if (err != cudaSuccess) {20        fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));21        return 1;22    }23    return 0;24}

<<<>>> 里到底填什么

<<<gridDim, blockDim>>> 这个三尖括号是 CUDA 独有的语法，叫执行配置（execution configuration）。第一个参数是 grid 里有多少个 block，第二个是每个 block 里有多少个线程。两者都可以是 dim3 三维类型，写成整数时等价于只用了 .x 维。

执行配置的几种写法
1kernel<<<2, 4>>>();                  // grid=(2,1,1), block=(4,1,1)2 3dim3 grid(16, 16);                   // grid=(16,16,1)4dim3 block(32, 8);                   // block=(32,8,1) = 256 线程5kernel<<<grid, block>>>();6 7// 完整形态还有两个可选参数：8//   第三个：每个 block 额外申请的动态共享内存字节数9//   第四个：把这次 launch 放进哪个 stream10kernel<<<grid, block, 48 * 1024, stream>>>();

编译与运行

用 nvcc 编译
1# 先确认驱动和卡2nvidia-smi3 4# -arch 指定目标架构，务必和你的卡对上，否则会走 JIT 甚至直接跑不了5#   sm_70 Volta / V100      sm_80 Ampere / A1006#   sm_75 Turing / T4       sm_86 Ampere / RTX 30xx7#   sm_89 Ada / RTX 40xx    sm_90 Hopper / H1008nvcc -O3 -arch=sm_80 -o hello 01_hello.cu9 10./hello

输出的顺序每次可能都不一样，比如 block 1 的线程先于 block 0 打印。这不是 bug，而是 CUDA 最重要的语义之一：block 之间没有任何执行顺序保证。GPU 可以按任意顺序、在任意 SM 上、以任意并发度调度它们。你写的 kernel 必须在这个前提下依然正确。

自测

先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。