在人工智能、科学计算和图形处理领域，CUDA（Compute Unified Device Architecture）已经成为加速计算的事实标准。从训练大语言模型到模拟物理现象，GPU凭借数千个计算核心，通过并行执行CUDA内核（Kernel）来大幅缩短计算时间。然而，当你写下 kernel<<<grid, block>>>(args); 并按下回车时，背后究竟发生了怎样精密而高效的硬件与软件协同？本文将为你拆解一个CUDA内核从CPU端启动到GPU完成计算的完整旅程。
第一步：CPU端的“发令枪”
CUDA程序通常以宿主端（CPU）代码为主。当你调用一个内核函数，编译器（如nvcc）已将其编译成PTX（并行线程执行）中间代码或直接生成SASS（流处理器汇编）指令。运行时，CPU线程执行 cudaLaunchKernel 这一底层API，它包含了内核函数指针、网格和线程块的维度配置、参数列表以及流（Stream）信息。此时，CPU不会等待GPU完成，而是立即将启动命令压入CUDA驱动维护的命令缓冲区，然后继续执行后续代码——这种异步设计是GPU高吞吐量的基石。
第二步：驱动层的“翻译与分发”
CUDA驱动收到启动命令后，首先检查当前设备（GPU）是否空闲或已有排队的任务。然后，驱动将内核的二进制代码从主机内存传输到GPU的全局内存（Global Memory）中的固定区域，并设置一个称为“控制单元”的结构，记录网格大小、共享内存大小等元数据。这些信息随后通过PCIe总线或NVLink互联发送到GPU的硬件前端：GigaThread引擎。
第三步：GigaThread引擎的“网格分拆”
GigaThread引擎是GPU上的任务调度中枢。它把整个网格（Grid）拆解为若干个线程块（Thread Block），并按照资源消耗（如寄存器用量、共享内存需求）将这些线程块分配到不同的流多处理器（SM）上。每个SM都有一个线程块调度器，它维护一个待分配的线程块队列。当一个SM的硬件资源（如计算单元、寄存器文件、共享内存）足够容纳一个新线程块时，调度器便从队列中取出一个线程块并激活它。
第四步：SM内部的“Warp调度与执行”
每个SM进一步将线程块内的线程划分为更小的执行单位——Warp（典型为32线程）。SM内的warp调度器会持续选择一条有就绪指令的warp发射到计算单元上。由于不同warp可能因内存访问延迟、分支分歧等原因处于不同状态，调度器采用零开销切换策略：当某个warp因等待内存读取而暂停时，立即切换到另一个可执行的warp，从而隐藏延迟。这种机制使得GPU即使拥有数千个线程，也能保持近乎满负荷的计算吞吐。
第五步：内存访问的“分层博弈”
内核执行过程中，线程可访问多层次内存：寄存器（最快，每线程独有）、共享内存（L1缓存与SM内共享，低延迟）、全局内存（大容量但高延迟）、纹理与常量缓存等。编译器会尽力将频繁使用的变量存入寄存器或共享内存，但当线程试图读取全局内存时，会触发缓存缺失（cache miss），访存请求必须经过L2缓存甚至直接访问显存（DRAM）。现代GPU通过合并访问（coalesced access）优化——当同一warp中线程访问连续地址时，硬件将多个请求合并为一次较大传输，大幅提升带宽利用率。
第六步：同步与屏障
内核内部可调用 __syncthreads() 进行线程块内同步，确保所有线程都已执行到同一位置后再继续。硬件通过barrier寄存器实现：每个线程到达屏障后递减计数器，直到计数器归零才放行。跨线程块的同步则必须通过内核结束或使用全局原子操作（如atomicAdd）来实现。这些机制保证了数据依赖正确性，但也可能引入性能瓶颈。
第七步：结果回传与清理
当所有线程块执行完毕，SM向GigaThread引擎报告完成。最终，一个中断信号通过PCIe通知CPU内核执行结束。此时，CPU端如果之前调用了 cudaDeviceSynchronize()，则会阻塞等待；或者可以通过 cudaStreamSynchronize(stream) 等待特定流完成。之后，调用 cudaMemcpy 将GPU显存中的计算结果传回主机内存，完成一次完整的CUDA内核启动周期。
结语
从CPU端的一条函数调用，到GPU内部千核齐发的并行狂欢，CUDA内核的执行是一场精密的交响乐：驱动调度、硬件分拆、warp发射、内存优化、同步退场，每个环节都经过数十年迭代优化。理解这一过程，不仅有助于写出更高效的GPU代码，也让我们窥见现代计算架构如何将“简单”的并行化愿景，转化为每秒亿亿次浮点运算的现实。随着CUDA生态向更多领域扩展，深入掌握“内核运行背后”的机制，将是每一个高性能计算从业者的必修课。