1. 课程导入与CUDA编程全景

在数值计算领域，CPU 长期承担着串行逻辑与复杂控制流的处理任务，而 GPU 凭借数千个流处理器核心，在大规模并行浮点运算上展现出数量级的吞吐优势。CUDA 正是 NVIDIA 提供的通用并行计算平台，它让开发者可以用接近 C/C++ 的语法，把计算任务映射到 GPU 的并行执行单元上。本课程的目标，是帮助大家从线程模型出发，逐步掌握编写、调试和优化数值核函数的能力。

要理解 CUDA 的定位，先要建立异构计算的基本图景。一个典型的 CUDA 程序运行在“主机 + 设备”的异构系统中：主机（Host）指 CPU 及其内存，设备（Device）指 GPU 及其显存。二者拥有各自独立的地址空间，通过 PCIe 总线连接。CPU 擅长分支判断、文件读写、任务调度；GPU 擅长对海量数据执行相同的算术操作。因此，一个 CUDA 程序的典型分工是：CPU 负责准备数据、分配显存、启动核函数、回收结果；GPU 负责在核函数中完成密集的并行计算。

上图展示了主机与设备之间的数据流与执行分工。可以看到，数据必须先通过显存拷贝进入设备，核函数才能读取；计算完成后，结果再拷回主机。这一“分配—拷贝—启动—回拷”的骨架，是几乎所有 CUDA 程序的共同结构。

接下来是 CUDA 线程模型，这是本课程最核心的概念。CUDA 把并行执行体组织成三级层次：线程（Thread）是最小执行单元；若干线程组成一个线程块（Block）；若干线程块组成一个网格（Grid）。每个线程通过内置变量 threadIdx、blockIdx、blockDim、gridDim 确定自己的全局编号。以最常见的向量加法为例，若数组长度为 N，每个线程处理一个元素，则全局索引可写为：

int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
    C[idx] = A[idx] + B[idx];
}

这里的 blockIdx.x * blockDim.x 定位到当前线程块在网格中的起始偏移，threadIdx.x 是线程在块内的偏移，两者相加得到全局下标。边界判断 idx < N 不可省略，因为线程总数通常向上取整，会略多于实际元素个数。理解这个索引公式，就理解了 CUDA 线程模型与数据元素之间的映射关系，后续的矩阵运算、归约、卷积都建立在这一映射之上。

开发环境方面，需要准备三样东西：支持 CUDA 的 NVIDIA 显卡、对应版本的显卡驱动，以及 CUDA Toolkit。安装完成后，用 nvcc --version 确认编译器可用，用 nvidia-smi 查看显卡型号与驱动版本。编写第一个程序时，源文件以 .cu 为扩展名，核函数用 __global__ 修饰，主机调用代码用 cudaMalloc 分配显存、cudaMemcpy 传输数据、cudaDeviceSynchronize 等待设备完成。编译命令形如：

nvcc -o vecadd vecadd.cu
./vecadd

一个值得强调的实践细节是错误检查。CUDA 的许多调用失败时不会中断程序，而是返回错误码。建议在每次显存分配、拷贝和核函数启动后，用 cudaGetLastError() 或封装宏检查返回值，否则很容易出现“结果全为零”却找不到原因的困境。例如核函数启动后立即调用 cudaGetLastError()，若返回非 cudaSuccess，说明网格或块维度配置有误。

从学习路径看，本课程将沿着“线程模型 → 核函数编写 → 内存层次 → 性能分析 → 典型数值算法”的主线推进。第一阶段的重点不是追求极致性能，而是建立正确的映射思维：数据如何分布到线程、线程如何协作、主机与设备如何配合。当你能独立写出一个通过边界检查、结果正确的向量加法核函数，并能解释每个线程在做什么，就已经跨过了 CUDA 编程最重要的门槛。后续的共享内存、同步、归约与矩阵分块，都是在这一基础模型上的自然延伸。