📚 课程目录第1章课程导入

课程导入：学习目标、课程结构、GPU计算基础与CUDA概述

1.1 学习目标与课程定位

本课程聚焦于英伟达GPU运行时（Runtime）体系中的内核调度与执行机制，目标不是停留在CUDA编程接口的调用层面，而是深入理解从CPU发起内核启动到GPU硬件完成计算这一完整链路中，运行时库、驱动、硬件三者如何协同。完成本章学习后，你应当能够：

描述GPU异构计算的基本执行模型，区分主机端（Host）与设备端（Device）的职责边界；
说出CUDA运行时（libcudart）与驱动API（libcuda）在架构层次上的关系；
理解内核（Kernel）启动时“网格-线程块-线程”三级层次如何映射到SM（流式多处理器）硬件资源；
建立后续章节（内存模型、流与并发、内核调度器、硬件队列等）的知识地图。

课程共30章，按“基础机制 → 运行时核心 → 调度深入 → 性能优化 → 排障实战”五个阶段递进。本章作为起点，先搭建GPU计算的整体框架，后续章节将逐一拆解每个环节的运行时实现细节。

1.2 GPU计算基础：为什么需要异构架构

CPU与GPU的设计哲学存在本质差异。CPU追求低延迟，拥有大容量缓存和复杂分支预测单元，擅长执行串行逻辑；GPU追求高吞吐，将大量晶体管用于浮点运算单元（FP32/FP64/Tensor Core），通过数千个线程的并行掩盖访存延迟。下表对比了两者在关键维度上的差异：

维度	CPU	GPU
核心数量	8~64个高性能核心	数千个轻量级CUDA核心
缓存策略	大容量L2/L3，复杂预取	较小L2，依赖线程切换隐藏延迟
执行模型	MIMD（多指令多数据）	SIMT（单指令多线程）
典型延迟	访存延迟约100周期	访存延迟约400~800周期，但可并行掩盖
适用场景	逻辑控制、串行任务	大规模并行数值计算、矩阵运算

CUDA（Compute Unified Device Architecture）是英伟达提出的通用并行计算平台。程序员编写在GPU上运行的函数称为内核（Kernel），通过特殊语法标记并由运行时库负责加载与调度。一个典型的CUDA程序结构如下：

// 主机端代码
__global__ void vectorAdd(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() {
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, n * sizeof(float));
    cudaMalloc(&d_b, n * sizeof(float));
    cudaMalloc(&d_c, n * sizeof(float));
    cudaMemcpy(d_a, h_a, n * sizeof(float), cudaMemcpyHostToDevice);
    // 内核启动：配置网格与线程块
    vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
    cudaMemcpy(h_c, d_c, n * sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

上述代码中，<<<gridSize, blockSize>>>语法是运行时库提供的启动接口，它最终会触发一系列复杂的内部操作——这正是本课程后续要深入剖析的核心内容。

关键认知：CUDA的“运行时”并非单一组件，而是包含用户态运行时库（cudart）、驱动API（cuda）以及内核驱动模块（nvidia.ko）的完整栈。用户代码通过cudart发起调用，cudart再与驱动交互，最终由驱动将工作提交给GPU前端。

1.3 内核执行模型：网格、线程块与线程

CUDA采用三级线程层次。内核启动时，程序员需要指定网格（Grid）中线程块（Block）的布局，以及每个线程块内线程（Thread）的数量。这种层次结构并非随意设计，它直接映射到GPU硬件的调度单元：

线程（Thread）：最小的执行单元，每个线程拥有独立的程序计数器和寄存器状态；
线程块（Thread Block）：由多个线程组成，块内线程可通过共享内存（Shared Memory）通信，并可进行同步（__syncthreads()）；
网格（Grid）：由多个线程块组成，块之间无法直接通信，只能通过全局内存间接交互。

硬件层面，GPU由多个流式多处理器（SM）构成。线程块被分配到SM上执行，一个SM可同时驻留多个线程块。线程块内的线程进一步划分为线程束（Warp）——通常32个线程为一束，这是硬件真正调度和执行的基本单位。下图展示了从软件层次到硬件映射的完整关系：

上图揭示了本课程的核心主线：软件线程层次与硬件SM资源之间的映射关系。运行时库和驱动的主要职责之一，就是高效地完成这种映射，并在资源受限时进行排队与调度。

1.4 CUDA运行时架构分层

CUDA软件栈从用户代码到硬件可划分为四个层次，理解每一层的职责对后续深入内核调度至关重要：

应用层：用户编写的CUDA C/C++代码，包含内核函数和主机端逻辑；
CUDA运行时库（libcudart）：提供cudaMalloc、cudaMemcpy、<<<>>>等高层API，管理设备内存、流、事件和模块加载；
CUDA驱动API（libcuda）：更底层的接口，直接与内核驱动通信，负责上下文管理、模块加载、内核参数绑定和启动提交；
GPU内核驱动与硬件：驱动将命令推送到GPU前端（Front-End），由GPU的工作分配器（Work Distributor）将线程块分发到各SM。

运行时库与驱动API之间存在明确的封装关系。运行时库内部会调用驱动API，但用户通常无需直接感知这一层。当用户调用kernel<<<grid, block>>>(args)时，实际发生的调用链如下：

用户代码
  → cudart: cudaLaunchKernel()   // 运行时库入口
    → 检查上下文与模块句柄
    → 配置内核参数（通过CUDAoccupancy或用户指定）
    → 调用驱动API: cuLaunchKernel()
      → 驱动验证参数合法性
      → 将内核启动信息封装为命令包
      → 写入GPU前端命令队列（Push Buffer）
        → GPU硬件调度器解析命令
        → 工作分配器将线程块分发至SM

  重要概念：内核启动是异步的。CPU发起cudaLaunchKernel后不会等待GPU执行完成，而是立即返回。真正的同步点由cudaDeviceSynchronize()或cudaMemcpy等操作触发。这种异步机制是GPU实现计算与数据传输重叠的基础，也是后续“流与并发”章节的核心前提。

1.5 常见误区与学习路径建议

初学者在理解GPU运行时机制时，常陷入以下误区，本课程将逐一澄清：

误区一：认为<<<grid, block>>>只是简单的语法糖，忽略了其背后涉及的参数合法性检查、内核属性查询、资源占用计算等运行时开销。
误区二：混淆“线程块调度到SM”与“线程束在SM内执行”两个层次。前者由硬件工作分配器完成，后者由SM内部的线程束调度器完成，两者处于不同的硬件流水线阶段。
误区三：认为所有线程块会同时执行。实际上，当SM资源不足时，部分线程块需要等待，直到前一批块完成并释放资源。
误区四：忽略运行时库与驱动API的版本兼容性。不同CUDA版本对内核启动参数布局、默认流语义（特别是cudaStreamLegacy与cudaStreamPerThread）的处理存在差异。

本课程后续将按以下路径深入：第2~5章剖析内存模型与数据传输；第6~9章讲解流、事件与并发执行；第10~15章深入内核启动的运行时内部实现，包括cudaLaunchKernel的完整流程、参数缓冲区布局、模块加载机制；第16~20章聚焦GPU硬件调度器、工作分配器与SM级调度；第21~25章讨论性能分析工具（Nsight Compute）与内核调优；第26~30章通过实际案例进行故障排查与深度复盘。