1. CUDA并行计算工程导论

在传统CPU程序里，我们习惯把任务写成一条顺序执行的指令流：取数、计算、写回，循环往复。CPU为这种模式做了极致优化——大容量缓存、分支预测、乱序执行，让单条指令流的延迟尽可能低。但当问题规模膨胀到数亿个像素、数百万个粒子或上千亿次矩阵乘加时，CPU的几十个核心就显得捉襟见肘。GPU换了一条路：它不追求单个线程跑得多快，而是用成千上万个轻量线程同时开工，用吞吐量换延迟。这就是异构计算范式的核心——CPU负责逻辑控制与串行部分，GPU负责数据密集的并行部分，两者通过PCIe或NVLink协同。

理解CUDA，先要理解它的执行模型。CUDA程序由主机端（Host，即CPU）和设备端（Device，即GPU）代码组成。主机端负责分配显存、拷贝数据、启动核函数；设备端执行核函数（kernel）。核函数被成千上万个线程并行执行，这些线程按线程—线程块—网格三级组织：32个线程组成一个warp，是硬件调度的基本单位；若干warp组成一个线程块（block），块内线程可通过共享内存通信；若干块组成一个网格（grid），对应一次核函数启动。线程通过blockIdx、threadIdx等内置变量确定自己的全局编号，从而决定处理哪一部分数据。

CUDA生态远不止一个编译器。从底层的驱动API、运行时API，到cuBLAS、cuFFT、cuDNN、TensorRT等加速库，再到Nsight系列性能分析工具、NCCL多卡通信库、CUDA Graph、统一内存等运行时特性，构成了完整的工程栈。实际项目中，我们很少从零手写所有算子——矩阵乘用cuBLAS，卷积用cuDNN，通信靠NCCL，自己专注写那些库覆盖不到的自定义核函数，这才是工程化的正确姿势。

学习路线建议分四步走：第一步，搭好环境，跑通第一个向量加法核函数，理解Host/Device内存模型；第二步，掌握线程索引、内存层次（全局内存、共享内存、寄存器）、同步机制；第三步，学会用Nsight Compute和Nsight Systems定位瓶颈，理解占用率、内存带宽、指令吞吐等指标；第四步，进入工程实践，处理多流并发、动态并行、多GPU通信等真实场景。下面这张图展示了从主机端到设备端的完整数据流与执行链，是理解CUDA程序结构的骨架。

环境搭建是第一步，也是最容易卡住新手的地方。你需要三样东西：一块支持CUDA的NVIDIA GPU（计算能力≥5.0即可，现代工程建议≥7.0）、对应版本的显卡驱动、以及CUDA Toolkit。在Ubuntu 22.04上，推荐用官方runfile安装，避免apt源版本混乱：

# 查看GPU型号与驱动版本
nvidia-smi

# 下载并安装CUDA Toolkit 12.4（示例）
wget https://developer.download.nvidia.com/compute/cuda/12.4.0/local_installers/cuda_12.4.0_550.54.14_linux.run
sudo sh cuda_12.4.0_550.54.14_linux.run

# 配置环境变量
echo 'export PATH=/usr/local/cuda-12.4/bin:$PATH' >> ~/.bashrc
echo 'export LD_LIBRARY_PATH=/usr/local/cuda-12.4/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc
source ~/.bashrc

# 验证
nvcc --version

安装完成后，写一个最小可运行的向量加法程序来验证环境。这个例子虽然简单，却包含了CUDA编程的全部关键要素：显存分配、数据拷贝、核函数定义与启动、结果回传与校验。

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

__global__ void vecAdd(const float* A, const float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) C[i] = A[i] + B[i];
}

int main() {
    const int N = 1 << 20;          // 约100万个元素
    size_t bytes = N * sizeof(float);
    float *h_A = (float*)malloc(bytes);
    float *h_B = (float*)malloc(bytes);
    float *h_C = (float*)malloc(bytes);
    for (int i = 0; i < N; i++) { h_A[i] = 1.0f; h_B[i] = 2.0f; }

    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    cudaDeviceSynchronize();

    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
    printf("C[0]=%f, C[N-1]=%f\n", h_C[0], h_C[N-1]);

    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);
    return 0;
}

编译命令为 nvcc -o vecadd vecadd.cu。运行后若输出 C[0]=3.000000, C[N-1]=3.000000，说明环境正常。这里有几个值得注意的工程细节：blocksPerGrid 用向上取整公式计算，确保覆盖所有元素；核函数内用 if (i < N) 做边界保护，防止越界访问；cudaDeviceSynchronize() 确保核函数执行完毕后再回传数据。这些看似琐碎的检查，在实际工程中是避免段错误和结果错误的底线。

最后要建立的一个认知是：CUDA不是“把CPU代码翻译成GPU代码”的魔法。它要求你重新思考数据的划分方式、线程的映射关系、内存的访问模式。一个在CPU上跑得很快的算法，直接搬到GPU上可能因为线程发散、内存不连续访问而变得很慢。工程化的学习路径，就是不断在“写得出”和“跑得快”之间迭代，用工具量化瓶颈，用架构知识指导优化。本章后续内容将围绕这条主线逐步展开。