2、CUDA编程基础：CUDA安装与验证、CUDA编程模型、简单向量加法实现

2.1 先说几句题外话

我们做深度学习的人，跟CUDA打交道是躲不掉的。我刚开始接触CUDA时，总觉得这东西很深奥，又是线程又是块的，搞得挺玄乎。说白了，CUDA就是让显卡干活的一套工具。

你想想看，CPU有十几个核心就很厉害了。而GPU呢？成百上千个核心。CUDA就是教你怎么指挥这些核心，让它们帮你算矩阵、做卷积。我在项目中遇到过一位同事，CPU代码写得飞起，一上GPU就懵。为什么？他总想把CPU的思维搬到GPU上，结果跑得比CPU还慢。

2.2 CUDA环境安装与验证

先搞定环境。我建议你按这个顺序来：

1. 检查显卡 — 先看看你的NVIDIA显卡支不支持CUDA。
2. 安装驱动 — 驱动版本要跟CUDA版本匹配。
3. 安装CUDA Toolkit — 包含编译器、库和工具。
4. 验证安装 — 跑几个命令确认一切正常。

2.2.1 检查显卡

打开终端，直接敲：

nvidia-smi

你会看到类似这样的输出：

+-----------------------------------------------------------------------------+
| NVIDIA-SMI 545.23.08    Driver Version: 545.23.08    CUDA Version: 12.3     |
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|===============================+======================+======================|
|   0  NVIDIA GeForce RTX 4090  On   | 00000000:01:00.0  On |                  Off |
| 30%   45C    P8    25W / 450W |  12345MiB / 24564MiB |      0%      Default |
+-----------------------------------------------------------------------------+

重点看右上角的CUDA Version。这个版本号就是你系统支持的最高版本。嗯，这里要注意，如果你装的是CUDA 12.3的驱动，那你最多只能用CUDA 12.3的工具包。

2.2.2 安装CUDA Toolkit

去NVIDIA官网下载对应版本的CUDA Toolkit。我通常用runfile方式安装，因为可控性高。当然你也可以用deb包或rpm包。

以CUDA 12.3为例，安装命令大致是这样：

wget https://developer.download.nvidia.com/compute/cuda/12.3.0/local_installers/cuda_12.3.0_545.23.08_linux.run
sudo sh cuda_12.3.0_545.23.08_linux.run

安装过程中会提示你同意协议，选择安装组件。我一般只勾选CUDA Toolkit和Driver（如果驱动还没装）。

装完后记得配环境变量：

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

你可以把这俩加到~/.bashrc里，省得每次重开终端都要敲一遍。

2.2.3 验证安装

环境配好了，怎么知道CUDA能不能用？我个人习惯用这几个命令来验证：

# 查看CUDA版本
nvcc --version

# 编译一个简单的测试
cd /usr/local/cuda-12.3/samples/1_Utilities/deviceQuery
sudo make
./deviceQuery

nvcc --version会显示编译器版本号。deviceQuery这个例子会打印出显卡详细信息，包括CUDA核心数、显存大小、支持的计算能力等。

如果看到类似下面这样的输出，就说明安装成功了：

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 12.3, CUDA Runtime Version = 12.3, NumDevs = 1, Device0 = NVIDIA GeForce RTX 4090, Compute Capability = 8.9
Result = PASS

2.3 CUDA编程模型

理解了CUDA编程模型，你就理解了GPU加速的精髓。咱们从三个层面来讲：网格、线程块、线程。

概念	说明	我常用的比喻
网格（Grid）	一个核函数启动后，所有线程的总集合	整支军队
线程块（Block）	一组可以协作的线程，共享一块内存	一个连队，可以互通消息
线程（Thread）	执行的最小单元	一个士兵

调用核函数时，我们要指定网格大小和块大小：

kernel_name<<<grid_size, block_size>>>(arguments);

grid_size是线程块的数量，block_size是每个块里的线程数。这两个参数可以是一维、二维或三维的。

我个人做深度学习项目时，最喜欢把数据做成一维的，因为卷积和全连接层的本质就是大量向量运算。当然，图像处理时二维块也很常用。

2.3.1 层级关系

每个线程都有自己唯一的ID，通过threadIdx、blockIdx、blockDim来定位：

int tid = blockIdx.x * blockDim.x + threadIdx.x;

这个公式一定要记住。我在项目中遇到过有人把公式写反，结果数据全乱了。你想想看，每个线程都拿到了错误的索引，那运算结果能对吗？

2.4 简单向量加法实现

理论说完了，咱们上手写一个最简单的例子：两个向量的加法。这在深度学习里很常见，比如两个特征图逐元素相加。

2.4.1 CPU版本

先看看CPU怎么实现的：

#include <iostream>
#include <vector>

void vector_add_cpu(const float* a, const float* b, float* c, int n) {
    for (int i = 0; i < n; ++i) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    int n = 1 << 20; // 约100万个元素
    std::vector<float> a(n, 1.0f), b(n, 2.0f), c(n, 0.0f);
    vector_add_cpu(a.data(), b.data(), c.data(), n);
    std::cout << "c[0] = " << c[0] << std::endl;
    return 0;
}

这代码大家都会写，循环一下就行了。但数据量一大，CPU就扛不住了。

2.4.2 GPU版本

现在换成CUDA实现。我会把代码分成三部分：设备内存分配、核函数、数据回传。

#include <iostream>
#include <cuda_runtime.h>

__global__ void vector_add_gpu(const float* a, const float* b, float* c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

int main() {
    int n = 1 << 20;
    size_t bytes = n * sizeof(float);

    // 1. 在GPU上分配内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 2. 初始化主机数据
    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;
    }

    // 3. 把数据拷贝到GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 4. 启动核函数
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;
    vector_add_gpu<<<grid_size, block_size>>>(d_a, d_b, d_c, n);

    // 5. 把结果拷回CPU
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    std::cout << "c[0] = " << h_c[0] << std::endl;

    // 6. 清理
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    delete[] h_a;
    delete[] h_b;
    delete[] h_c;

    return 0;
}

这段代码里有几个要点：

• __global__ 修饰符表示这是一个核函数，在GPU上执行，从CPU调用。
• cudaMalloc 在GPU上分配内存，不能直接用new或malloc。
• cudaMemcpy 负责CPU和GPU之间的数据传输。方向参数要写对，我一开始经常搞混。
• if (tid < n) 这个边界检查不能少。因为总线程数可能超过元素个数，多余的线程就不干活。

2.4.3 编译与运行

用nvcc编译：

nvcc -o vector_add vector_add.cu
./vector_add

如果一切顺利，你会看到输出：

c[0] = 3

1+2=3，没错。

2.5 本章核心知识图

下面这张图展示了CUDA编程模型从Host到Device的完整流程：