课程导入：CUDA编程的定位、学习路径、环境搭建与第一个简单程序

CUDA编程的定位与核心原理

要理解CUDA的定位，首先需要明确CPU与GPU在架构设计上的根本差异。CPU（中央处理器）由少数几个高性能核心（通常4-16个）组成，每个核心拥有大容量缓存（L1/L2/L3）和复杂的控制单元，擅长处理分支密集、依赖性强、延迟敏感的任务。而GPU（图形处理器）则由数千个简单计算核心组成，这些核心被组织成流多处理器（Streaming Multiprocessor, SM），每个SM内部包含多个CUDA Core、共享内存（Shared Memory）和寄存器文件。GPU的设计哲学是"以吞吐量换延迟"——单个线程的执行速度并不快，但通过大规模线程并行，总计算吞吐量远超CPU。

CUDA编程模型的核心是异构计算：主机（Host，即CPU及其内存）负责串行逻辑和任务调度，设备（Device，即GPU及其显存）负责数据并行计算。程序员需要编写两个部分的代码：在CPU上运行的普通C/C++代码，以及在GPU上运行的核函数（Kernel Function）。核函数使用__global__关键字声明，通过特殊的语法kernelName<<<gridDim, blockDim>>>(args)启动。当核函数被启动时，GPU会创建大量线程，这些线程按照线程层次结构组织：最上层是网格（Grid），网格由多个线程块（Block）组成，每个线程块包含多个线程（Thread）。这种层次结构直接映射到GPU硬件：线程块被调度到SM上执行，块内的线程可以通过共享内存和同步机制协作。

学习路径规划与知识体系

CUDA编程的学习路径应当遵循"由浅入深、理论结合实践"的原则。根据多年教学经验，建议将学习过程划分为五个阶段，每个阶段都有明确的目标和验证标准。

在学习过程中，建议遵循"每学一个概念，立即编写验证程序"的原则。CUDA编程是实践性极强的技能，只看书不写代码等于没学。同时，要养成使用性能分析工具的习惯——从第一阶段就开始接触Nsight Systems和Nsight Compute，理解程序的运行时间和瓶颈所在，而不是等到优化阶段才学习。

开发环境搭建：从驱动到工具链

搭建CUDA开发环境是学习的第一步，也是许多初学者容易遇到问题的环节。一个完整的CUDA开发环境包含三个层次：NVIDIA显卡驱动、CUDA Toolkit、开发工具（IDE和调试分析器）。下面详细介绍每个组件的安装与验证方法。

1. NVIDIA驱动安装

NVIDIA驱动是操作系统与GPU硬件之间的桥梁，它提供了CUDA运行时所需的底层支持。安装驱动前，首先确认你的GPU型号是否支持CUDA。可以通过以下命令查询：

# Linux系统
lspci | grep -i nvidia
# Windows系统（PowerShell）
wmic path win32_VideoController get name

确认GPU型号后，推荐从NVIDIA官网下载对应型号的最新驱动。对于Linux系统，建议使用发行版自带的驱动管理工具（如Ubuntu的ubuntu-drivers）或NVIDIA官方.run文件安装。安装完成后，运行nvidia-smi命令验证驱动是否正常工作，该命令会显示GPU型号、驱动版本、显存使用情况等信息。

2. CUDA Toolkit安装

CUDA Toolkit包含编译器（nvcc）、CUDA运行时库、数学库（cuBLAS、cuFFT等）以及开发工具。安装方式根据操作系统不同有所差异：

# Ubuntu/Debian（使用官方APT仓库）
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb
sudo dpkg -i cuda-keyring_1.1-1_all.deb
sudo apt-get update
sudo apt-get install cuda

# Windows（使用安装包）
# 从NVIDIA官网下载cuda_*.exe安装包，选择自定义安装，建议勾选CUDA和Nsight工具

安装完成后，需要配置环境变量。在Linux的~/.bashrc中添加：

export PATH=/usr/local/cuda/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH

验证安装是否成功：

nvcc --version
# 输出应包含Cuda compilation tools, release 12.x, V12.x

3. Nsight工具与IDE配置

Nsight是NVIDIA提供的专业开发工具套件，包含Nsight Systems（系统级性能分析）和Nsight Compute（内核级性能分析）。这些工具对于后续的性能优化至关重要。推荐使用Visual Studio Code作为IDE，安装以下扩展：

• CUDA C++：提供语法高亮和代码补全
• Nsight Visual Studio Code Edition：集成调试和性能分析
• C/C++：基础语言支持

在VS Code中配置c_cpp_properties.json，将CUDA头文件路径添加到includePath中：

{
  "configurations": [
    {
      "name": "Linux",
      "includePath": ["${workspaceFolder}/**", "/usr/local/cuda/include"],
      "defines": [],
      "compilerPath": "/usr/local/cuda/bin/nvcc",
      "cStandard": "c17",
      "cppStandard": "c++17",
      "intelliSenseMode": "linux-gcc-x64"
    }
  ],
  "version": 4
}

第一个CUDA程序：向量加法实战

让我们通过经典的向量加法程序，完整走一遍CUDA开发的流程。这个案例虽然简单，但包含了CUDA编程的所有核心要素：内存分配、数据拷贝、核函数编写与启动、结果验证。

案例背景：假设我们需要计算两个长度为N=1024的浮点数组A和B的元素和，结果存入数组C。在CPU上，这需要循环N次；在GPU上，我们可以创建N个线程，每个线程计算一个元素的和。

完整代码：

#include <cstdio>
#include <cuda_runtime.h>

// 核函数：每个线程计算一个元素
__global__ void vectorAdd(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() {
    int n = 1024;
    size_t bytes = n * sizeof(float);

    // 1. 分配主机内存并初始化
    float *h_A = new float[n];
    float *h_B = new float[n];
    float *h_C = new float[n];
    for (int i = 0; i < n; i++) {
        h_A[i] = 1.0f;
        h_B[i] = 2.0f;
    }

    // 2. 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);

    // 3. 将数据从主机拷贝到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

    // 4. 配置核函数启动参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;

    // 5. 启动核函数
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);

    // 6. 将结果拷贝回主机
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

    // 7. 验证结果
    bool correct = true;
    for (int i = 0; i < n; i++) {
        if (h_C[i] != 3.0f) {
            correct = false;
            break;
        }
    }
    printf("Result: %s\n", correct ? "PASS" : "FAIL");

    // 8. 释放内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    delete[] h_A;
    delete[] h_B;
    delete[] h_C;

    return 0;
}

编译与运行：

nvcc -o vector_add vector_add.cu
./vector_add
# 输出：Result: PASS

代码拆解与关键点分析：

常见误区与调试技巧

初学者在编写CUDA程序时，经常会遇到以下几类问题。提前了解这些陷阱，可以节省大量调试时间。

误区一：忽略错误检查。CUDA API函数（如cudaMalloc、cudaMemcpy）返回错误码，核函数启动也可能失败（如配置参数错误）。建议封装一个错误检查宏：

#define CHECK_CUDA(call) do { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \
                cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
} while(0)

误区二：主机和设备内存混淆。CPU代码不能直接访问设备指针（d_A），GPU核函数也不能直接访问主机指针（h_A）。所有数据传递必须通过cudaMemcpy显式完成。一个常见的错误是在核函数中访问主机指针，导致段错误或未定义行为。

误区三：线程索引计算错误。当数组长度不是线程块大小的整数倍时，需要向上取整计算网格大小，并在核函数内做边界检查。例如，N=1000，block=256时，blocksPerGrid = (1000+255)/256 = 4，最后一个块只有232个有效线程，必须通过if (i < n)过滤。

误区四：忽视内存带宽限制。向量加法是内存密集型（Memory-bound）任务，其性能受限于PCIe带宽和显存带宽，而非计算能力。在后续章节中，我们会学习如何通过共享内存、合并访问等技术优化访存模式。

调试工具推荐：除了printf（核函数内支持printf，但输出顺序不保证），建议使用Nsight Systems查看程序时间线，确认数据传输和核函数执行的时间占比；使用Nsight Compute查看内核的占用率、内存吞吐量等指标。对于复杂错误，可以使用cuda-gdb（Linux）或Nsight调试器（Windows）进行断点调试。