第四章：CUDA 编程基础：GPU 架构简介、CUDA 编程模型、核函数编写与调用

各位同学，欢迎来到第四章。说实话，这一章是整个域控制器实战课程的「硬骨头」，也是你从嵌入式工程师跨入高性能计算领域的门槛。我当年第一次接触 CUDA 时，看着满屏的 __global__ 和 <<>>，心里直犯嘀咕：这玩意儿跟单片机上的中断服务程序到底有啥区别？

别急，咱们一步步来。先搞清楚 GPU 到底是个什么怪物，再谈怎么让它干活。

4.1 GPU 架构简介：从「单核大力士」到「万核蚂蚁雄兵」

CPU 和 GPU 的设计哲学完全不同。CPU 是「单核大力士」，擅长处理复杂的逻辑分支和串行任务。GPU 则是「万核蚂蚁雄兵」，它拥有成千上万个简单计算核心，专门用来处理数据并行任务。

我打个比方。CPU 就像一位顶尖外科医生，能做精细的心脏搭桥手术。GPU 就像一万个建筑工人，每人搬一块砖，合力盖起摩天大楼。你让外科医生去搬砖，浪费人才；你让建筑工人去做手术，那得砸锅。

在汽车电子领域，典型的 GPU 并行任务包括：

  • 图像处理：摄像头输入的每一帧图像，每个像素点的滤波、边缘检测都可以独立计算。
  • 矩阵运算：传感器融合中的卡尔曼滤波、深度学习推理中的卷积操作。
  • 点云处理：激光雷达的每个点云数据，做坐标变换或聚类。

英伟达的 GPU 架构经历了从 Fermi、Kepler、Maxwell、Pascal、Volta、Turing 到 Ampere 的演进。咱们域控制器上常用的 Orin 芯片，用的是 Ampere 架构。它的核心计算单元叫 SM（Streaming Multiprocessor，流式多处理器）。

每个 SM 内部包含：

  • CUDA Cores：整数和浮点运算单元，Orin 的每个 SM 有 128 个。
  • Tensor Cores：专门为深度学习矩阵乘法设计的加速单元，做 D=A*B+C 这种操作时，吞吐量是 CUDA Core 的几十倍。
  • Shared Memory（共享内存）：同一 Block 内的线程可以访问的快速缓存，延迟极低。
  • Register File（寄存器文件）：每个线程私有的最快存储。

核心概念：GPU 的并行能力来自于「大量线程 + 快速上下文切换」。当某个线程在等待内存访问时，硬件会立刻切换到另一个就绪的线程，几乎零开销。这就是 GPU 能隐藏内存延迟的秘密武器。

我记得第一次在 Orin 上做性能分析时，发现某个 kernel 的 occupancy（占用率）只有 30%。说白了就是 SM 上同时运行的线程太少，硬件没法通过切换来隐藏延迟。后来调整了 block 大小和寄存器使用量，occupancy 提到 80%，性能直接翻倍。嗯，这里要注意，不是 occupancy 越高越好，但太低肯定有问题。

4.2 CUDA 编程模型：主机与设备的「主仆关系」

CUDA 编程模型把系统分为两个角色：

  • Host（主机）：CPU 及其内存（系统内存）。负责控制逻辑、数据准备、启动 kernel。
  • Device（设备）：GPU 及其显存（全局内存）。负责执行大规模并行计算。

你想想看，这就像项目经理（CPU）和施工队（GPU）的关系。项目经理制定计划、准备材料，然后一声令下：「干活！」施工队就闷头执行。项目经理不会去搬砖，施工队也不会去画图纸。

典型的 CUDA 程序流程如下：

  1. 分配显存：在 GPU 上分配全局内存。
  2. 数据拷贝：将数据从系统内存拷贝到显存。
  3. 启动 Kernel：调用 GPU 上的核函数。
  4. 结果回传：将计算结果从显存拷贝回系统内存。
  5. 释放资源：释放显存。

这里有个坑，我踩过不止一次。数据拷贝是通过 PCIe 总线进行的，带宽远低于显存带宽。如果你在循环里频繁做 cudaMemcpy，性能会惨不忍睹。我曾经接手过一个项目，同事把每帧图像都单独拷贝一次，结果帧率只有 5fps。改成异步拷贝 + 双缓冲后，直接跑到 30fps。

避坑指南：尽量减少 Host 与 Device 之间的数据传输。能一次传完就别分多次。如果必须频繁传输，考虑使用 cudaMemcpyAsync 结合流（Stream）来重叠计算与拷贝。

4.3 核函数编写与调用：让 GPU 动起来

核函数（Kernel）就是在 GPU 上执行的函数。它的定义方式很特别，需要用 __global__ 修饰符，并且返回类型必须是 void。

来看一个最简单的例子：两个数组相加。

// 核函数定义
__global__ void vectorAdd(const float* A, const float* B, float* C, int N) {
    // 计算当前线程的全局索引
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查，防止越界
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

// 主机端调用
int main() {
    int N = 1 << 20;  // 1048576 个元素
    size_t bytes = N * sizeof(float);
    
    // 分配主机内存
    float *h_A, *h_B, *h_C;
    h_A = (float*)malloc(bytes);
    h_B = (float*)malloc(bytes);
    h_C = (float*)malloc(bytes);
    
    // 初始化数据（略）
    
    // 分配设备显存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);
    
    // 拷贝数据到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
    
    // 配置执行配置
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    
    // 启动核函数
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // 拷贝结果回主机
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
    
    // 释放资源
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);
    
    return 0;
}

这里有几个关键点，我展开说说：

4.3.1 线程索引计算

blockIdx.x 是当前线程块在网格中的索引，blockDim.x 是每个线程块包含的线程数，threadIdx.x 是当前线程在块内的索引。三者组合起来，就能得到全局唯一的线程 ID。

为什么这么设计？因为 GPU 的线程是分层次组织的。网格（Grid）包含多个线程块（Block），每个 Block 包含多个线程。这种层次结构是为了让同一个 Block 内的线程可以共享数据（通过共享内存），并且可以同步。

4.3.2 执行配置

<<<blocksPerGrid, threadsPerBlock>>> 这个语法，我第一次看到时觉得像外星符号。它的含义是：用 blocksPerGrid 个 Block，每个 Block 有 threadsPerBlock 个线程，总共 blocksPerGrid * threadsPerBlock 个线程并行执行。

我个人习惯把 threadsPerBlock 设为 128 或 256，这是大多数 GPU 的「甜点值」。设得太小（比如 32），SM 的利用率不够；设得太大（比如 1024），每个线程分到的寄存器太少，反而会 spill 到本地内存，拖慢速度。

4.3.3 边界检查

注意核函数里的 if (idx < N)。为什么需要这个？因为 blocksPerGrid 是向上取整算的，实际启动的线程数可能略多于元素个数。如果不做边界检查，多余的线程会访问越界内存，轻则数据错误，重则程序崩溃。

个人经验：在调试阶段，我习惯在核函数开头加一个 printf 打印线程 ID，确认索引计算是否正确。虽然 printf 在 GPU 上很慢，但调试时用用无妨。正式发布前记得删掉。

4.4 实战中的注意事项

讲完了基础，咱们聊聊在汽车电子域控制器上写 CUDA 程序时，我遇到过的几个典型问题。

问题 现象 原因 解决方案
核函数不执行 程序无报错，但结果不对 忘记检查 cudaGetLastError() 每次 kernel 启动后调用 cudaGetLastError() 和 cudaDeviceSynchronize()
性能瓶颈 GPU 利用率低 数据传输与计算未重叠 使用 CUDA Stream 实现异步操作
内存泄漏 程序运行一段时间后崩溃 cudaMalloc 后忘记 cudaFree 使用 RAII 封装，或定期检查显存使用量
Bank Conflict 共享内存访问变慢 多个线程同时访问同一 Bank 调整数据布局，使用 padding 技术

我曾经在调试一个车道线检测算法时，发现核函数计算结果总是有随机噪声。排查了两天，最后发现是 cudaMalloc 分配的显存没有初始化，而核函数里又假设初始值为 0。从那以后，我养成了习惯：分配显存后立即用 cudaMemset 清零，或者用 cudaMallocManaged 统一内存管理。

4.5 本章小结

这一章的内容，说白了就是三件事：

  • GPU 架构：记住 SM、CUDA Core、Shared Memory 这三个概念。它们是性能优化的基石。
  • 编程模型：Host 负责控制，Device 负责计算。尽量减少数据传输。
  • 核函数：__global__ 定义，<<<grid, block>>> 启动，注意索引计算和边界检查。

下一章我们会深入共享内存和同步机制，那是真正拉开性能差距的地方。到时候我会拿一个实际的传感器融合案例，手把手带大家优化。今天就到这里，各位回去把上面的 vectorAdd 例子跑通，有任何问题随时交流。

← 上一章 📖 返回目录 下一章 →