CUDA 编程入门：从 GPU 架构到核函数优化

作者：新兰2026.11.04 12:09浏览量：0

简介：本文面向希望系统理解 CUDA 编程原理的开发者，从 GPU 多核架构、线程组织、内存模型等基础概念出发，逐步讲解核函数编写、执行配置与性能优化方法。读者将掌握 CUDA 程序的基本结构、线程束调度机制、共享内存与 L2 缓存优化思路，并能完成简单的向量加法与矩阵乘法示例验证，为后续深入学习高性能计算打下基础。

一、教程目标

本教程旨在帮助读者建立对 CUDA 编程基本原理的系统认知，并完成从零到一的入门实践。具体而言，学完本文后你应当能够：

  1. 理解 GPU 与 CPU 在架构层面的核心差异，以及这种差异如何影响编程模型；
  2. 掌握 CUDA 编程中线程、线程块、网格、线程束等核心概念及其层级关系；
  3. 编写并运行一个完整的 CUDA 核函数，完成向量加法等基础并行计算任务；
  4. 理解全局内存、共享内存、L2 缓存等存储层次对性能的影响；
  5. 建立基本的性能分析意识，知道从哪些维度去判断和优化核函数效率。

本教程适合有一定 C/C++ 基础，但对 GPU 并行编程接触较少的开发者。如果你已经写过一些 CUDA 代码，但对其底层执行机制理解不够透彻，也可以通过本文重新梳理知识体系。

二、GPU 架构基础：从多核到线程束

理解 CUDA 编程的第一步，是理解 GPU 的硬件组织方式。与 CPU 相比，GPU 的核数量多得多，但单个核的控制逻辑相对简单。现代高端 GPU 通常包含上百个流式多处理器（Streaming Multiprocessor，简称 SM），例如某些旗舰级数据中心 GPU 拥有 130 个以上的 SM。

每个 SM 内部包含以下关键资源：

  • 计算单元：包括专门用于矩阵乘法的张量核心（Tensor Core），以及用于一般标量运算的算术逻辑单元（ALU），后者支持 INT32、FP32、FP64 等不同精度的运算。张量核心主要服务于神经网络中的矩阵乘累加操作，而通用 ALU 则处理更广泛的标量计算。
  • 调度器：每个 SM 内有多个线程束调度器（Warp Scheduler），负责将指令分发给计算单元执行。
  • 寄存器文件与共享内存：存储层次中最靠近计算单元的部分，访问速度最快，但容量有限。
  • L1 缓存：与共享内存通常共用同一块物理存储，用于缓存全局内存数据。

从宏观上看，所有 SM 都通过片上互连访问 L2 缓存，L2 缓存再连接到片外的全局内存（DRAM）。L2 缓存的容量通常远小于全局内存，但访问延迟低得多。因此，提升 L2 缓存命中率是 CUDA 性能优化的关键目标之一，这与 CPU 程序中的缓存优化思路有相似之处，只是 GPU 场景下这种优化对最终性能的影响往往更加显著。

三、CUDA 线程模型与执行层级

CUDA 的编程模型建立在多层级的线程组织方式之上。理解这个层级结构，是写好核函数的前提。

3.1 网格、线程块与线程

CUDA 程序中的并行计算由大量线程共同完成。这些线程按照以下结构组织：

  • 线程（Thread）：执行的最小单位，每个线程有唯一的局部标识。
  • 线程块（Thread Block）：一组线程的集合，被分配到同一个 SM 上执行。同一线程块内的线程可以通过共享内存进行数据交换，也可以进行块内同步。
  • 网格（Grid）：一个核函数启动时创建的所有线程块的集合。整个网格对应一次完整的核函数调用。

这种层级设计使得 CUDA 程序天然具备可扩展性：当 GPU 拥有更多 SM 时，更多的线程块可以被并行调度，同一个核函数无需修改即可在更大规模的 GPU 上运行。

3.2 线程束：指令分发的真正单位

虽然代码层面我们以线程为单位编写逻辑，但硬件实际以线程束（Warp）为单位执行指令。一个线程束由 32 个连续的线程组成，它们在同一时刻执行相同的指令，只是操作的数据不同。这种执行模式属于单指令多数据（SIMD）架构的一种变体，在 GPU 中被广泛采用。

线程束的存在带来两个重要启示：

第一，分支问题需要格外谨慎。 当同一个线程束内的 32 个线程因为条件判断走上不同的执行路径时，硬件会串行执行各条路径，导致性能下降。这被称为线程束发散（Warp Divergence）。尽量避免线程束内部的严重分支，是核函数优化的重要技巧之一。

第二，内存访问模式对性能影响巨大。 当线程束内所有线程访问连续且对齐的内存地址时，硬件可以将多个访问请求合并为较少的传输事务，这被称为内存合并访问（Memory Coalescing）。如果访问模式分散，则会产生大量额外的内存事务，显著拖慢执行速度。

3.3 线程束调度器的作用

线程束调度器是 SM 内部负责指令分发的硬件单元。它的任务是从已就绪的线程束中选择一个，将其下一条指令发送到相应的计算单元。由于单个线程束在执行长延迟操作（如全局内存访问）时会阻塞，SM 需要同时维护大量活跃线程束，以便在某个线程束等待时快速切换到另一个线程束继续执行。这种机制被称为延迟隐藏（Latency Hiding）。

这意味着，即使每个线程束自身的执行速度受限于内存延迟，只要 SM 上有足够多的线程束可供切换，整体吞吐量仍然可以保持高位。因此，在资源允许的情况下，尽量为核函数配置更多的线程块和线程数，通常有利于性能表现。

四、核函数编写入门

4.1 核函数的基本形态

核函数（Kernel）是在 GPU 上执行的函数，使用 __global__ 修饰符声明。一个最简单的向量加法核函数可以这样写：

  1. __global__ void vectorAdd(const float* a, const float* b, float* c, int n) {
  2.     int idx = blockIdx.x * blockDim.x + threadIdx.x;
  3.     if (idx < n) {
  4.         c[idx] = a[idx] + b[idx];
  5.     }
  6. }

这里需要理解几个关键内建变量：

  • blockIdx.x：当前线程块在网格中的索引；
  • blockDim.x：每个线程块包含的线程数；
  • threadIdx.x：当前线程在其线程块内的索引。

三者组合起来，就可以为每个线程计算出一个全局唯一的索引 idx。边界判断 idx < n 是必要的，因为线程总数可能大于数据元素数量，需要防止越界访问。

4.2 执行配置

核函数的调用使用三尖括号语法指定执行配置。例如：

  1. int threadsPerBlock = 256;
  2. int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
  3. vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(a, b, c, n);

第一个参数是网格中的线程块数量，第二个参数是每个线程块内的线程数。线程块大小通常选择 32 的倍数（如 128、256、512），原因是线程束大小为 32，选择 32 的倍数可以避免线程束内部出现未使用的线程槽位。每个 SM 能容纳的线程块数量和线程总数有硬件上限，这一点需要在选择执行配置时综合考虑。

4.3 结果验证

核函数执行完成后，需要将结果从 GPU 显存复制回主机内存，然后进行校验。通用流程如下：

  1. 在主机端准备输入数据；
  2. 在 GPU 端分配对应的显存空间；
  3. 将输入数据从主机复制到 GPU；
  4. 启动核函数；
  5. 将输出结果从 GPU 复制回主机；
  6. 释放在 GPU 上分配的资源；
  7. 在主机端验证结果的正确性。

对于向量加法示例，可以逐元素比较 GPU 计算结果与 CPU 串行计算结果，若所有元素差值在容差范围内，即可确认核函数逻辑正确。

五、存储层次与性能优化方向

5.1 全局内存与访问模式

全局内存位于 GPU 芯片之外的 DRAM 上，是所有线程都可以访问的存储空间，但访问延迟最高。优化全局内存访问的核心原则是内存合并访问：让同一个线程束内的 32 个线程在可行的情况下访问一段连续的内存地址。

例如，在线程束内，如果线程 t 访问数组的第 t 个元素，且数组地址按 128 字节对齐，则这 32 次访问可以被合并为一次或少数几次内存传输事务。如果线程访问的地址是跳跃式的（如跨步长为 2 或更大），则会产生多倍的内存事务，降低有效带宽利用率。

5.2 共享内存的意义

每个 SM 上都有一定容量的共享内存，由该 SM 上运行的所有线程块共享。共享内存的访问延迟远低于全局内存，适合用于存放线程块内需要多次重复读取的数据。常见的优化模式是将数据从全局内存加载到共享内存，然后在计算过程中从共享内存读取，从而减少对全局内存的依赖。

使用共享内存时需要注意两点：一是共享内存容量有限，需要合理规划数据规模；二是在写入和读取之间需要调用 __syncthreads() 进行块内同步，确保所有线程都完成写入后再开始读取。

5.3 L2 缓存的作用

L2 缓存位于所有 SM 与全局内存之间，是所有 SM 共享的片上存储。当多个线程块重复访问相同的数据区域时，L2 缓存能够显著降低延迟和带宽压力。CUDA 中的某些机制允许开发者对 L2 缓存的使用进行策略性控制，例如通过访问策略窗口设置数据在 L2 中的保留优先级。对于需要反复遍历数据集的算法，合理利用 L2 缓存可以带来明显的性能提升。

六、常见问题与排查思路

6.1 核函数运行结果不正确

这是入门阶段最常见的问题。排查顺序建议如下：

  1. 首先检查索引计算是否正确。blockIdx、blockDim、threadIdx 的组合是否准确映射到了目标数据位置；
  2. 检查是否存在越界访问。边界判断条件是否覆盖了所有可能出现的情况；
  3. 确认内存分配和释放操作是否匹配；
  4. 确认核函数调用后的错误检查代码是否执行。调用核函数后应立即调用错误检查函数（如 cudaGetLastError()），而不是等到程序末尾才检查。

6.2 程序运行报内存相关错误

这类问题通常与内存分配失败、非法访问或未初始化指针有关。需要检查以下几点：

  • 主机端和 GPU 端的指针是否在使用前都已正确分配；
  • 显存分配的大小是否与数据实际占用的字节数一致；
  • 复制操作中指定的字节数是否正确；
  • 核函数内部是否有未初始化变量的使用。

6.3 性能表现低于预期

如果在功能正确的前提下性能不理想，可以从以下几个方面分析：

  • 线程块大小是否合适：过小的线程块无法充分占用 SM 资源，过大的线程块可能超出硬件限制；
  • 内存访问是否合并：访问模式是否导致大量未合并的内存事务；
  • 活跃线程束数量是否充足：通过分析工具检查 SM 占用率，确认是否有足够的线程束用于隐藏延迟；
  • 是否存在线程束发散：核函数内部是否有条件分支导致同一线程束内线程走不同路径。

七、深入学习的建议路径

在掌握了上述基础概念之后，建议按照以下顺序进一步深化：

  1. 实践经典并行算法：从向量加法开始，逐步尝试矩阵乘法、归约求和、前缀和等算法。矩阵乘法是理解共享内存、线程协作和张量核心使用方式的良好切入点。
  2. 分析核函数性能指标：学习使用 GPU 厂商提供的性能分析工具，关注内存吞吐量、计算利用率、占用率等关键指标，将感性认知转化为量化分析。
  3. 研究张量核心编程：对于面向深度学习推理和训练的开发者，理解张量核心的编程模型和数据布局要求非常关键。
  4. 了解 CUDA 流与异步执行：当单个核函数的优化空间被充分挖掘后，可以进一步学习如何通过多流并行来提升整体吞吐量。
  5. 探索更高层级的编程抽象：在实际工程项目中，直接编写底层 CUDA 代码的维护成本较高。了解行业常见的并行编程抽象与库，可以在保证性能的同时显著提升开发效率。

八、总结

CUDA 编程的入门重点不在于记忆 API，而在于理解 GPU 的硬件执行模型与编程模型之间的映射关系。核心要点可以归纳为：

  • GPU 由多个 SM 构成，每个 SM 内部包含多种计算单元和线程束调度器；
  • CUDA 程序以网格、线程块、线程的层级结构组织并行计算；
  • 线程束是硬件执行指令的基本单位，分支控制和内存访问模式对性能至关重要；
  • 存储层次从寄存器、共享内存到 L2 缓存、全局内存，访问延迟逐级递增；
  • 性能优化的核心思路是提高计算利用率、减少不必要的内存访问、提升内存合并度和缓存命中率。

掌握了这些基础概念，你就已经具备了进一步学习高性能 GPU 计算的必要框架。接下来，动手实践是最好的老师——先写一个最简单的核函数并让它正确运行，再逐步增加复杂度，在调试和优化的过程中加深对每一个概念的理解。

最热文章