如何学习 CUDA 编程？一份面向开发者的系统指南

不大耳朵图图
2 阅读8分钟

在大模型、科学计算、图形处理席卷一切的今天，GPU 不再只是"显卡"，而是算力的主战场。而驾驭 GPU 的钥匙，正是 CUDA。

这篇文章写给所有想入门 CUDA 编程的开发者——无论你是做 AI 基础设施、高性能计算，还是单纯想弄懂"GPU 到底是怎么并行的"。我会从学习路径、环境搭建、第一个 Kernel、核心概念、优化思路到参考资料，给你一条清晰可执行的路线。

一、先想清楚：你为什么要学 CUDA？

CUDA 是 NVIDIA 推出的并行计算平台和编程模型，它让开发者可以用 C++、Python、Fortran 等语言调用 GPU 的大规模并行算力，在深度学习、科学计算和高性能计算（HPC）领域被广泛采用。

但 GPU 不是"万能加速器"。一个常见的误区是：

💡 GPU 只在计算密集、高并行度的任务上才比 CPU 快；串行逻辑、分支复杂的代码，放到 GPU 上反而更慢。

所以在开始之前，先确认你的场景是否属于以下几类：

  • AI 训练 / 推理优化：自定义算子、TensorRT 插件、推理引擎开发
  • 科学计算 / HPC：矩阵运算、流体仿真、分子动力学
  • 图像处理 / 计算机视觉：滤波、变换、特征提取的大规模并行
  • 个人兴趣：想真正搞懂 GPU 的执行模型和内存层次

如果只是用 PyTorch 跑模型，你不需要写 CUDA；但如果你想做上面的"底层工作"，CUDA 就是绕不过的坎。

二、学习前的知识储备

CUDA C++ 是 ANSI C 的扩展，所以你需要：

  1. 扎实的 C++ 基础：指针、引用、内存管理、模板

  2. 基本的并行计算思维：理解"将一个大问题拆成上千个小问题同时求解"

  3. 一块 NVIDIA GPU（核心要求）：

    • 入门级：RTX 3060 / 4060 性价比不错
    • 没有本地 GPU？Google Colab 提供免费的 Tesla T4，AutoDL、阿里云 GPU 实例也是选项
    • AMD / Intel GPU 不支持 CUDA，它们有各自的 ROCm / oneAPI

三、推荐的学习路径

参考 NVIDIA 官方编程指南的结构，以及社区的共识，我为你梳理了一条 6 阶段路线：

C/C++ 基础
   ↓
并行计算思维
   ↓
CUDA 基础（线程模型、内存模型）
   ↓
经典算法实现（向量加法、矩阵乘、规约）
   ↓
性能优化（内存、占用率、Stream）
   ↓
实战项目（AI 推理 / 图像处理 / 科学计算）
   ↓
高级主题（多 GPU、NVLink、Tensor Core）

四、搭建开发环境

1. 安装 CUDA Toolkit

CUDA Toolkit 包含了 GPU 加速库、调试和优化工具、C++ 编译器（nvcc）以及运行时库。

Linux（Ubuntu）下验证安装：

# 查看 GPU 状态
nvidia-smi

# 查看 CUDA 编译器版本
nvcc --version

# 如果未安装 nvcc
apt install nvidia-cuda-toolkit
nvcc --version

Windows​ 下推荐 Visual Studio + CUDA Toolkit 的组合，或者用 VS Code + Nsight 的轻量方案。

📌 没有本地 GPU 的同学：可以用 ssh连接到有 GPU 的远程 Linux 服务器开发，或者用 Docker 容器挂载 GPU 资源。

2. 你的第一个 CUDA 程序

新建 hello_world.cu：

#include <stdio.h>

__global__ void helloFromGPU() {
    printf("Hello, World from GPU! thread=%d\n", threadIdx.x);
}

int main() {
    printf("Hello from CPU!\n");
    
    // 启动 1 个 Block，每个 Block 5 个 Thread
    helloFromGPU<<<1, 5>>>();
    
    // 核函数调用是异步的，需要同步等待
    cudaDeviceSynchronize();
    return 0;
}

编译运行：

nvcc hello_world.cu -o hello_world
./hello_world

输出：

Hello from CPU!
Hello, World from GPU! thread=0
Hello, World from GPU! thread=1
...
Hello, World from GPU! thread=4

这段代码里有三个 CUDA 核心概念：

  • __global__ ：函数限定符，表示这个函数在 host 调用、在 device 执行，且只能返回 void
  • <<<1, 5>>> ：核函数启动语法，第一个数字是 Grid 中 Block 的数量，第二个是每 Block 的 Thread 数
  • cudaDeviceSynchronize() ：因为核函数调用是异步的，CPU 不会等 GPU 算完，所以需要显式同步

五、CUDA 编程模型的核心心智

1. 异构计算：Host vs Device

CUDA 是一个异构环境——CPU（Host）和其内存，与 GPU（Device）和其内存，在硬件和软件上都是隔离的。数据要在两者之间显式搬运。

一个典型的 CUDA 程序遵循 5 步结构：

  1. 在 GPU 上分配内存（cudaMalloc）
  2. 从 CPU 内存拷贝数据到 GPU（cudaMemcpy）
  3. 调用核函数执行计算
  4. 将结果从 GPU 拷贝回 CPU
  5. 释放​ GPU 内存（cudaFree）

2. 线程层次结构

CUDA 的线程组织是分层的：

Grid（网格）
  └── Block（线程块，可被调度到 SM 上）
        └── Thread（线程，执行最小单位）
  • Grid​ → 一次核函数启动的所有 Block
  • Block​ → 一组线程，同一个 Block 内的线程可以共享共享内存、可以同步
  • Thread​ → 单个执行单元

核函数启动时的配置 <<<grid, block>>>就是在设定这个层次。

3. 内存层次结构

决定 CUDA 程序性能的两个模型：内存层次结构和线程层次结构。

  • 全局内存（Global Memory） ：容量大但延迟高，Host 和 Device 之间主要在这个层面搬运数据
  • 共享内存（Shared Memory） ：Block 内线程共享，速度远快于全局内存，是性能优化的关键
  • 寄存器（Register） ：每个线程私有，速度最快

六、动手写一个向量加法

理解了模型，我们来写一个真正有意义的程序——向量加法：

#include <stdio.h>
#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() {
    const int N = 1024;
    size_t bytes = N * sizeof(float);

    // 1. Host 端数据
    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] = i * 1.0f;
        h_B[i] = i * 2.0f;
    }

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

    // 3. 数据拷贝 Host -> Device
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

    // 4. 启动核函数：256 个线程/Block，需要多少 Block 就多少
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 5. 结果拷回 Host
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

    // 验证结果
    printf("C[0]=%f, C[100]=%f, C[1023]=%f\n", h_C[0], h_C[100], h_C[1023]);

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

    return 0;
}

这段代码值得你反复琢磨的地方：

  • blockIdx.x * blockDim.x + threadIdx.x是每个线程的全局索引计算方式
  • (N + threadsPerBlock - 1) / threadsPerBlock是经典的向上取整计算 Block 数的写法
  • if (i < N)是必须的边界保护——因为线程总数通常是 threadsPerBlock的整数倍，可能大于 N

七、从"能跑"到"跑得快"：性能优化方向

写出正确的 CUDA 程序只是开始，真正的功夫在优化。以下几个方向是必经之路：

1. 合并内存访问（Coalesced Access）

当一个 Warp 中的 32 个线程访问连续的内存地址时，GPU 可以将这些访问合并成极少的内存事务，吞吐量大幅提升；反之，跨步访问（Strided Access）会让性能断崖式下跌。

⚠️ 矩阵转置就是一个典型的两难：合并读必然导致离散写，合并写必然导致离散读。解决方案是用共享内存做中转。

2. 善用共享内存

共享内存的速度远高于全局内存。经典的优化模式是：

  1. 将全局内存的数据加载到共享内存
  2. 在共享内存中完成计算
  3. 将结果写回全局内存

这样可以大幅减少对全局内存的访问次数。

3. 占用率（Occupancy）

占用率指每个 SM 上活跃线程束占最大可执行线程束的比例。合理的 Block 大小、寄存器使用量、共享内存用量都会影响占用率。NVIDIA 提供了 Occupancy Calculator 工具辅助计算。

4. 异步执行与 Stream

CUDA 的核函数调用、内存拷贝默认都是异步的。通过 Stream​ 机制可以让多个核函数、内存拷贝重叠执行，隐藏延迟，最大化 GPU 利用率。

5. 使用 Nsight 工具链

NVIDIA Nsight 是一套强大的开发、调试、性能分析工具集。养成用 Nsight Systems​ 做系统级 profiling、Nsight Compute​ 做 kernel 级分析的习惯，比盲目猜测瓶颈高效 10 倍。

八、实战项目建议

光学不练假把式。当你掌握了基础后，强烈建议做以下项目：

  1. 手写矩阵乘法 vs cuBLAS：感受"自己写的"和"工业级库"的差距，理解为什么优化这么难
  2. GPU 版高斯模糊：图像处理的经典案例，涉及共享内存 tiling 优化
  3. 规约（Reduction）算法：从朴素实现到使用 warp shuffle 指令的极致优化
  4. 纯 CUDA 实现简单 CNN 前向传播：不依赖 PyTorch，理解深度学习框架的底层
  5. 多 GPU 编程：用 NVLink 或多卡通信库做 GPU 间协作

九、常见误区

❌ "GPU 比 CPU 快，所以所有任务都该用 GPU"

错。串行、分支密集的任务 GPU 反而更慢，选对工具才是关键。

❌ "PyTorch 就够了，没必要学底层 CUDA"

应用层确实不用。但要做推理引擎优化、自定义算子、HPC 研究，CUDA 是基本功。

❌ "显存越大越好，核心数不重要"

显存决定能放多大模型，核心数决定速度，两者缺一不可。

❌ 忽视核函数的异步性

新手最常犯的错：核函数启动后就立刻读结果，忘记 cudaDeviceSynchronize()，导致读到垃圾数据。

avatar
文章
阅读
粉丝