第二章：CUDA编程基础——从零搭建你的GPU开发环境

说实话，很多朋友刚开始学CUDA，第一关就卡在了环境搭建上。我见过有人折腾了三天，结果发现是驱动版本不对。嗯，今天咱们就把这事一次性说清楚。

2.1 驱动、工具包、编译器——三件套缺一不可

CUDA开发环境说白了就三样东西：NVIDIA驱动、CUDA Toolkit、编译器。我习惯把它们比作「引擎、工具箱和扳手」。

2.1.1 NVIDIA驱动

这是最底层的。你的GPU要干活，得先有驱动。我个人建议直接去NVIDIA官网下载最新版，别用系统自带的。为什么？因为系统自带的驱动往往版本太老，后面装CUDA Toolkit时会报错。

⚠️ 我曾经踩过的坑： 有一次我在Ubuntu上直接用apt安装驱动，结果CUDA 11.8死活装不上。查了半天，发现驱动版本是470，而CUDA 11.8需要驱动版本≥520。所以，先装驱动，再装Toolkit，顺序别搞反。

2.1.2 CUDA Toolkit

这个工具包包含了CUDA编译器（nvcc）、各种库（如cuBLAS、cuFFT）、以及调试工具。说白了，你写好的.cu文件，就是靠nvcc编译成GPU能执行的代码。

下载时注意：版本号要匹配。比如你的显卡是RTX 3090，那CUDA 11.x以上都支持。但如果是老卡（比如GTX 750），可能只能用CUDA 10.x。你想想看，这就像给自行车装火箭发动机——不是不行，是没必要。

2.1.3 编译器

Windows上我推荐用Visual Studio（2019或2022），Linux上直接用g++就行。nvcc本质上是个「包装器」，它会把你的.cu文件拆成两部分：主机代码（CPU部分）交给g++或cl.exe编译，设备代码（GPU部分）自己编译。

💡 小技巧： 安装完Toolkit后，在命令行输入 nvcc --version，如果能看到版本号，说明环境基本OK了。

2.2 第一个CUDA程序——向量加法

学编程语言，第一个程序通常是Hello World。但在CUDA里，我习惯用向量加法作为入门。为什么？因为它简单、直观，而且能清楚看到CPU和GPU的协作流程。

来看代码：

#include <stdio.h>

// GPU内核函数：两个向量相加
__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    int N = 1024;
    size_t size = N * sizeof(float);

    // 1. 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);

    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_A[i] = i * 1.0f;
        h_B[i] = i * 2.0f;
    }

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

    // 3. 把数据从CPU拷贝到GPU
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 4. 启动内核（1个Block，1024个Thread）
    vecAdd<<<1, 1024>>>(d_A, d_B, d_C, N);

    // 5. 把结果从GPU拷回CPU
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }

    // 6. 释放内存
    free(h_A); free(h_B); free(h_C);
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);

    return 0;
}

这段代码里，__global__ 关键字告诉编译器：这个函数要在GPU上运行。<<<1, 1024>>> 是CUDA特有的语法，表示启动1个Block，每个Block里有1024个Thread。

🔧 避坑指南： 我曾经在写第一个程序时，忘了加 cudaMemcpyDeviceToHost 这一步，结果打印出来的全是0。你想想看，数据还在GPU里没拷回来，CPU当然拿不到结果。所以，记得检查数据流向。

2.3 线程层次结构——Grid、Block、Thread

CUDA的线程模型，说白了就是三级组织架构：

  • Thread（线程）：最基础的计算单元，每个线程执行相同的代码（但处理不同的数据）。
  • Block（线程块）：一组线程的集合，它们可以共享内存（Shared Memory），也能通过同步指令（__syncthreads()）协同工作。
  • Grid（网格）：一组Block的集合，一个Grid对应一个内核函数（Kernel）的启动。

我习惯用「军队」来比喻：Thread是士兵，Block是班，Grid是连。每个士兵做自己的事，但同一个班的士兵可以互相配合。

来看一张图，帮你理清关系：

Grid（网格） Block (0,0) Block (1,0) T0 T1 T2 T3 T4 T5 T0 T1 T2 T3 T4 T5 一个Grid包含多个Block，每个Block包含多个Thread

在代码里，我们通过 blockIdx.x 和 threadIdx.x 来定位当前线程。比如上面向量加法的例子：

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

这个公式的意思是：当前线程的全局ID = 当前Block的ID × Block大小 + 线程在Block内的ID。说白了，就是给每个线程一个独一无二的编号，好让它知道自己该处理哪个数据。

2.4 内存模型简介——CPU和GPU的数据桥梁

CUDA的内存模型，我刚开始学的时候也觉得头大。其实你只要记住一句话：CPU不能直接访问GPU内存，GPU也不能直接访问CPU内存。数据必须通过PCIe总线来搬运。

来看这张表，一目了然：

内存类型 位置 访问权限 速度 生命周期
Global Memory（全局内存） GPU显存 所有Thread可读写 慢（约200-400 GB/s） 程序结束
Shared Memory（共享内存） GPU芯片内 同一Block内Thread共享 快（约1-2 TB/s） Block结束
Register（寄存器） GPU核心内 单个Thread私有 最快 Thread结束
Local Memory（局部内存） GPU显存（但逻辑上私有） 单个Thread私有 慢（同Global） Thread结束
Constant Memory（常量内存） GPU显存（带缓存） 所有Thread只读 快（有缓存） 程序结束
💡 实战经验： 我在做矩阵乘法优化时，发现把频繁访问的数据放到Shared Memory里，性能直接提升了5倍。你想想看，Global Memory就像去图书馆借书，Shared Memory就像把书放在自己桌上——哪个快？不言而喻。

最后，记住CUDA编程的六步流程：

  1. 分配主机内存（malloc）
  2. 分配设备内存（cudaMalloc）
  3. 把数据从CPU拷贝到GPU（cudaMemcpy Host→Device）
  4. 启动内核（Kernel Launch）
  5. 把结果从GPU拷回CPU（cudaMemcpy Device→Host）
  6. 释放内存（free + cudaFree）

嗯，这一章的内容就到这里。环境搭好了，第一个程序跑通了，线程模型也理解了——你已经迈出了GPU编程的第一步。后面的路还长，但方向对了，就不怕远。

专注资料整理
← 上一章 📖 返回目录 下一章 →