一、CUDA入门：GPU架构概述、CUDA编程模型、CUDA环境搭建与工具链

各位同学，欢迎来到《CUDA与TensorRT边缘计算部署实战》的第一章。

说实话，每次讲GPU架构，我都想起自己刚入行那会儿。那时候我对着NVIDIA的架构白皮书，看了三天三夜，愣是没搞明白SM和SP到底是个啥关系。后来在项目中调优一个图像处理算法，才真正体会到——不理解硬件，写出来的CUDA代码就是瞎蒙。

所以这一章，咱们先把地基打牢。我会用最直白的话，把GPU的底裤给你扒干净。

1.1 GPU架构概述：从“很多核”到“很多线程”

CPU和GPU最大的区别是什么？

CPU是“精英小团队”——几个大核，每个核都特别强，能处理各种复杂任务。GPU是“民工大部队”——成百上千个小核，每个核能力一般，但人多力量大，适合干重复性的体力活。

你想想看，渲染一张4K图片，需要计算几百万个像素的颜色值。每个像素的计算逻辑几乎一样，只是数据不同。这种场景，GPU就是天生的王者。

1.1.1 关键硬件概念

我们先搞清楚几个核心概念。这些词你以后会天天见。

术语	全称	通俗理解
SM	Streaming Multiprocessor	GPU里的“计算单元组”，相当于一个车间
SP / CUDA Core	Streaming Processor	SM里的“计算工人”，真正干活的
Warp	线程束	32个线程组成一队，一起执行指令
Global Memory	全局显存	GPU的大仓库，所有线程都能访问，但慢
Shared Memory	共享内存	SM内部的小仓库，快，但容量小

嗯，这里要注意：Warp是GPU调度的最小单位。什么意思？就是你写代码时觉得是“一个线程一个线程在执行”，但硬件层面其实是32个线程绑在一起跑的。这个特性会直接影响你的代码性能，后面我们会反复提到。

1.1.2 我踩过的坑：架构理解不到位

我曾经在一个项目中，用GPU做视频编解码加速。一开始我天真地以为，只要把循环改成并行就完事了。结果跑起来发现，性能还不如CPU。

后来一分析，问题出在“线程发散”上。因为我的算法里有大量if-else分支，同一个Warp里的32个线程，有的走if，有的走else。硬件没办法，只能先执行if的线程，再执行else的线程。32个线程的活儿，硬是干了两遍。

所以记住：尽量让同一个Warp里的线程走相同的执行路径。这是GPU编程的第一条铁律。

1.2 CUDA编程模型：你只需要记住三个层次

CUDA的编程模型，说白了就是三个层次：Grid、Block、Thread。

我习惯这么理解：

• Grid：整个任务，比如“处理一张图片”
• Block：任务的一个子集，比如“处理图片的左上角区域”
• Thread：真正干活的单元，比如“处理一个像素”

每个Block会被分配到某个SM上执行。一个SM可以同时处理多个Block，但Block之间是独立的，不能互相通信（除非用全局显存，但那样很慢）。

1.2.1 第一个CUDA程序：向量加法

咱们直接上代码。这是CUDA的“Hello World”——向量加法。

// 文件名：vector_add.cu
#include <stdio.h>

// GPU核函数
__global__ void vecAdd(float *A, 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() {
    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. 把数据从主机拷贝到设备
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 4. 配置核函数执行参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    // 5. 启动核函数
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 6. 把结果从设备拷贝回主机
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

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

    // 7. 清理
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);

    return 0;
}

1.2.2 代码里藏着什么？

你看这个核函数：

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

这行代码是CUDA编程的灵魂。它把三维的线程索引（blockIdx, threadIdx）映射到一维的数据索引。为什么这么写？因为GPU的线程是三维组织的，但我们的数据通常是一维数组。

说白了，就是给每个线程分配一个“工号”，然后告诉它：“你负责处理第i个数据”。

1.3 CUDA环境搭建与工具链

光说不练假把式。咱们得把环境搭起来。

1.3.1 你需要什么？

• 硬件：一块NVIDIA显卡（GTX 1050以上都行，实在没有用云GPU）
• 驱动：NVIDIA Driver（去官网下载，版本别太老）
• CUDA Toolkit：包含nvcc编译器、CUDA库、调试工具
• cuDNN：深度学习加速库（后面TensorRT会用到）

1.3.2 安装步骤（Ubuntu 20.04为例）

我建议用Linux开发CUDA。Windows也能用，但坑比较多。

# 1. 检查显卡
lspci | grep -i nvidia

# 2. 安装驱动（推荐用官方runfile）
sudo apt purge nvidia-*
wget https://us.download.nvidia.com/XFree86/Linux-x86_64/550.54.14/NVIDIA-Linux-x86_64-550.54.14.run
sudo sh NVIDIA-Linux-x86_64-550.54.14.run

# 3. 安装CUDA Toolkit
wget https://developer.download.nvidia.com/compute/cuda/12.4.0/local_installers/cuda_12.4.0_550.54.14_linux.run
sudo sh cuda_12.4.0_550.54.14_linux.run

# 4. 配置环境变量
echo 'export PATH=/usr/local/cuda-12.4/bin:$PATH' >> ~/.bashrc
echo 'export LD_LIBRARY_PATH=/usr/local/cuda-12.4/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc
source ~/.bashrc

# 5. 验证安装
nvcc --version
nvidia-smi

1.3.3 工具链一览

工具	用途	我的评价
nvcc	CUDA编译器	核心工具，必须掌握
nvidia-smi	监控GPU状态	日常必备，看显存、温度、利用率
nsys	性能分析（Nsight Systems）	调优神器，能看CPU/GPU时间线
ncu	内核分析（Nsight Compute）	深入分析每个核函数的瓶颈
cuda-gdb	GPU调试器	调试核函数里的bug

我个人习惯是：先用 nvidia-smi 看一眼显存够不够，然后用 nsys 抓一下整体性能，最后用 ncu 定位具体问题。这套组合拳，基本能解决90%的性能问题。

1.4 本章小结

这一章我们干了三件事：

1. 搞懂了GPU的硬件架构——SM、Warp、显存层次
2. 学会了CUDA编程模型——Grid、Block、Thread三层结构
3. 搭好了开发环境——驱动、CUDA Toolkit、工具链

下一章，我们会深入CUDA的内存管理。到时候我会告诉你，为什么 cudaMalloc 和 cudaMemcpy 用不好，你的程序会慢得像蜗牛。

嗯，今天就到这儿。记得把环境搭好，把向量加法的代码跑通。有问题随时交流。

---

公众号:蓝海资料掘金营，微信deep3321