1. CUDA 简介

1. GPU

GPU (Graphics Processing Unit) 是一种专门设计用于快速操作和修改内存，以加速大规模数据并行计算的微处理器。

GPU 最初专为图像和图形相关运算设计，但因其高度并行的架构，现已演化成通用并行处理器。与 CPU 不同，GPU 将绝大部分晶体管用于计算单元（ALU），而非复杂的控制逻辑和缓存。

GPU 内部拥有数千个简单小核，适合执行数据并行、计算密集型任务。CPU 内部则是几个到几十个复杂大核，适合执行串行、逻辑复杂任务。

 

1. GPU 内部布局

  • 一颗 GPU 包含多组 GPC（Graphics Processing Cluster，图形处理集群）

  • 每个 GPC 内含 多个 TPC（Texture Processing Cluster，图纹理处理集群）

  • 每个 TPC 内含 1 个或 2 个 SM（流式处理器，Streaming Multiprocessor）

  • 每个 SM 内部包含：CUDA Core、Tensor Core（AI 加速）、RT Core（光线追踪）、寄存器文件、L1 缓存/共享内存、纹理单元（部分在 SM 内/旁）等

  • GPC 内还有光栅引擎（图形用，纯计算卡会被移除）

  • 芯片外围分布着 多个 L2 缓存分区，每分区耦合 一组 ROP 和 一或多个 32/64 位内存控制器，直接连接 GDDR/HBM 显存

  • 所有 SM/GPC 通过 片上交叉互联或 Mesh 网络 统一访问 L2 缓存和显存

  • 外部互联部分包括 PCIe / NVLink / NVSwitch 控制器

  • 固定功能单元如视频编解码器（NVDEC/NVENC）、显示引擎 等布置在芯片一侧（计算卡通常会移除它们）

一张 GPU 是由数十个 SM 集群、共享大容量 L2 缓存、高带宽显存接口、外部互联构成。而决定一切计算能力的，正是 SM 的内部结构

 

2. SM

流式多处理器（Streaming Multiprocessor，SM）是 NVIDIA GPU 架构中最核心的计算单元。所有 CUDA 线程最终都会被分配到一个个 SM 上执行，因此 SM 的设计直接决定了 GPU 的并行能力、能效和功能特性。现代 SM 内部通常被划分为4个独立的处理分区（Sub-cores），它们共享 L1 缓存与共享内存

 

SM 内部结构：

  • CUDA Core

    基本算术单元，执行单精度浮点 (FP32) 和整数 (INT) 运算。现代架构配备了独立的 FP32 和 INT32 数据路径，使得浮点计算和寻址控制等整型计算可以并发执行，互不阻塞

  • FP64 Core

    专用于高精度科学计算的算术单元。在计算卡（如 A100）中大量配置，而在游戏卡中为了节省面积通常只保留极少数量

  • Tensor Core

    张量核心，专为深度学习与 AI 加速设计的矩阵运算单元，在一个时钟周期内即可完成庞大的混合精度矩阵乘加（MAC）运算

  • RT Core

    光线追踪加速核心，负责快速计算光线与三角形求交、遍历包围盒层次结构 (BVH) 等

  • 特殊功能单元 (SFU)

    执行超越函数，如正弦、余弦、倒数平方根、对数等，延迟比 CUDA Core 高

  • 加载/存储单元 (LSU)

    负责计算内存地址，并向共享内存、L1 缓存或全局内存发出数据读写请求

  • Warp 调度器与分发单元

    SM 的控制中枢。每个分区拥有独立的调度器，负责在每个时钟周期挑选状态就绪的线程束（Warp），并将指令分发到对应的硬件执行单元

  • 寄存器文件

    SM 内部速度最快、带宽最高的存储。每个 SM 拥有海量寄存器（如容量达 256KB），用于保存所有驻留 Warp 的线程上下文，实现零开销的线程切换以掩盖延迟

  • 共享内存 / L1 缓存

    一块可配置的高速片上 SRAM。L1 缓存由硬件自动管理，而共享内存完全暴露给开发者，是同一个线程块（Block）内各个线程进行高速数据交换和协同计算的核心媒介

  • 常量缓存 / 纹理缓存

    针对具有特定访问模式的数据（如只读、具备二维空间局部性）设计的专用缓存，用于加速对常量内存和纹理内存的读取

 

并行执行模型：SIMT

GPU 使用单指令多线程 (SIMT) 架构，多个线程执行相同的指令但处理不同的数据，其核心在于层级调度与极速上下文切换

  • 线程块 (Thread Block)

    任务分配的基本单元。一个 Block 会完整分配到单个 SM 上，并在生命周期内独占为其分配的物理资源（寄存器与共享内存）。Block 内的线程可以共享该 SM 的共享内存，并进行轻量级的屏障同步

  • 线程束 (Warp)

    硬件调度的最小单位。SM 不会逐个管理线程，而是以 32 个线程为一组 的 Warp 作为调度和执行的基本单位

  • 调度策略

    Warp 调度器在每个时钟周期从池中挑选就绪（操作数已就位）的 Warp，并将其下一条指令发射给执行单元。Warp 内的 32 个线程在同一时刻执行同一条指令，但处理各自的数据

  • 分支发散 (Warp Divergence)

    若 Warp 内的线程因 if-else 分支走向不同路径，硬件必须将两条分支序列化执行。执行分支 A 时，走向分支 B 的线程会被硬件掩蔽（Masked）并空转等待，这会破坏 SIMT 的并行效率，造成算力浪费

  • 延迟隐藏 (Latency Hiding)

    这是 SM 最关键的能力。当一个 Warp 因访存或长延迟计算而暂停时，调度器会立刻切换到另一个就绪的 Warp 继续执行。只要有足够多的活跃 Warp，计算单元就能一直保持忙碌，规避内存延迟

 

3. 架构演进

GeForce 256 (1999)：首个 GPU，硬件支持变换与光照（T&L），固定功能管线

GeForce 3 (2001)：首次引入可编程顶点着色器，像素着色器可进行有限计算，nfiniteFX 引擎

GeForce FX (2003)：完全可编程的顶点和像素着色器（Shader Model 2.0+），CineFX 架构

GeForce 6 (2004)：Shader Model 3.0，支持动态分支和 64 位 HDR，首次引入 SLI 多卡技术

Tesla (2006)：统一着色器架构（流处理器可执行各类着色器），引入 CUDA，彻底结束固定管线

Fermi (2010)：首个完整的 GPGPU 架构，可配置 L1 缓存/共享内存，支持 C++ 与 ECC，真正通用计算

Kepler (2012)：引入动态并行（GPU 可自行发射新内核），Hyper‑Q 多任务，SMX 大幅扩大规模

Maxwell (2014)：效能比飞跃，SM 拆分为 4 个独立子核心，实现极致能效的控制逻辑

Pascal (2016)：NVLink 高速互联，HBM2 显存，统一内存，FP16 半精度双倍吞吐，计算预占

Volta (2017)：首个张量核心（Tensor Core），专为矩阵乘加设计，FP16 吞吐飞跃，独立线程调度

Turing (2018)：首次引入 RT Core 实现实时光线追踪，张量核心支持 INT8/INT4，网格着色器

Ampere (2020)：第三代张量核心，支持稀疏矩阵加速、TF32 精度，MIG 多实例 GPU 硬件分区

Hopper (2022)：Transformer 引擎，FP8 精度，支持动态重编译/异步执行，线程块集群，DPX 指令

Blackwell (2024)：第五代张量核心，引入 FP4 等新精度格式，更大规模 AI 算力，NVLink 5 互联

 

2. CUDA

CUDA（Compute Unified Device Architecture） 是 NVIDIA 推出的并行计算平台与编程模型。它允许开发者使用 C/C++、Fortran、Python 等语言直接控制 GPU 硬件进行通用计算，无需将问题伪装成图形渲染。

CUDA 的核心思想是异构计算：CPU（主机）负责串行与控制逻辑，GPU（设备）负责高度并行的数据计算。

 

CUDA 安装
下载安装操作系统和显卡对应版本的 CUDA Toolkit，使用以下 CMD 指令验证安装：

1
powershell nvcc --version

 

开发工具

Visual Studio Code：安装 NVIDIA Nsight Visual Studio Code Edition 拓展

Visual Studio：支持 CUDA 项目

 

调试与分析工具

Nsight Visual Studio（Windows）/ cuda-gdb（Linux）：CUDA 调试器

compute-sanitizer：内存访问越界、泄漏等检查工具

Nsight Systems：系统级性能分析（CPU‑GPU 交互、时间线）

Nsight Compute：CUDA kernel 级微架构分析

 

CMake 设置 CUDA

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
project(gpu LANGUAGES CXX CUDA)

file(GLOB_RECURSE GPU_SOURCES CONFIGURE_DEPENDS
    "${CMAKE_CURRENT_SOURCE_DIR}/*.cpp"
    "${CMAKE_CURRENT_SOURCE_DIR}/*.cu"
)

file(GLOB_RECURSE GPU_HEADERS CONFIGURE_DEPENDS
    "${CMAKE_CURRENT_SOURCE_DIR}/*.h"
    "${CMAKE_CURRENT_SOURCE_DIR}/*.cuh"
)

find_package(CUDAToolkit REQUIRED)
set(CMAKE_CUDA_STANDARD 20)
set(CMAKE_CUDA_STANDARD_REQUIRED ON)
set(CMAKE_CUDA_EXTENSIONS OFF)

add_library(${PROJECT_NAME} STATIC ${GPU_SOURCES} ${GPU_HEADERS})
target_include_directories(${PROJECT_NAME} PUBLIC
    "${CMAKE_CURRENT_SOURCE_DIR}/include"
)

target_link_libraries(${PROJECT_NAME} PRIVATE CUDA::cudart)

 

2. 编程模型

对于 C++，CUDA 提供两套 API：

  • Runtime API（cuda_runtime.h） ：高级、易用，隐藏了初始化、上下文管理等细节
  • Driver API（cuda.h）：低级、灵活，允许直接控制上下文和模块加载，适合框架开发者或需要细粒度控制的场景

两者可共存但同一内核启动需保持 API 统一

 

1. 主机与设备

Cuda 编程假设系统由主机（Host）和设备（Device）构成：

  • 主机：只 CPU 及内存，负责串行逻辑、数据准备和内核调用
  • 设备：GPU 及显存，执行高度并行的计算任务

开发着需要显式管理主机与设备之间的数据传输，并在设备上启动计算内核，典型的程序流程为：

  1. 分配主机内存并初始化数据
  2. 分配设备内存（cudaMalloc）
  3. 将数据从主机拷贝到设备（cudaMemcpy）
  4. 启动内核执行并行计算
  5. 将结果从设备拷回主机
  6. 释放设备内存

 

2. 核函数

核函数是在设备上执行的函数，使用 __global__ 声明，通过特定的 <<<...>>> 语法从主机端调用：

1
2
3
4
5
__global__ void vectorAdd(const float* A, const float* B, float* C, int N) 
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N) C[i] = A[i] + B[i];
}

调用方式：

1
vectorAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, N);

核函数在每个线程上独立执行，所有线程执行同一段代码，但通过不同的内置变量计算出各自的索引，从而操作不同的数据

 

3. 线程组织

在 CUDA 的编程模型中，线程被组织成：Grid（网格）、Block（线程块）、Thread（线程）三个层次：

  • Thread

    最基本的执行单元。每个线程执行相同的核函数，但处理不同的数据。每个线程都有自己的局部内存（Local Memory）和寄存器（Registers）

  • Block

    由一组相互协作的线程组成。同一个 Block 内的线程可以通过共享内存（Shared Memory）进行极速的数据交换，并且可以进行同步操作（__syncthreads()）

  • Grid

    由一组 Block 组成。一个 Kernel 函数的所有 Block 构成一个 Grid。不同 Block 之间的线程是相互独立的，它们不能直接同步，执行顺序也是未知的

 

CUDA 提供四个内置的 uint3 类型的变量（包含 xyz 三个维度）来让线程定位自身和要处理的数据：

内置变量 含义 作用域
gridDim 当前 Grid 的维度大小（即包含多少个 Block） Grid 全局可见
blockIdx 当前 Block 在 Grid 中的索引坐标 仅 Block 内可见
blockDim 当前 Block 的维度大小（即包含多少个 Thread） Block 全局可见
threadIdx 当前 Thread 在 Block 中的索引坐标 仅 Thread 内可见

在启动核函数时，我们通过 <<<grid, block>>> 语法来定义 gridDim 和 blockDim

线程组织的核心目的，就是计算出当前线程在整个 Grid 中的全局唯一索引，从而将其映射到数组或矩阵的具体位置上：

  • 一维组织，常用于处理一维数组

    1
    2
    3
    4
    5
    6
    7
    8
    9
    // 假设启动配置： <<< gridDim.x, blockDim.x >>>
    __global__ void addKernel(float* a, float* b, float* c) 
    {
        int tid = blockIdx.x * blockDim.x + threadIdx.x;
        if (tid < N) 
        {
            c[tid] = a[tid] + b[tid];
        }
    }

  • 二维组织，常用于处理图像处理或矩阵运算

    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    // 假设启动配置： <<< dim3(nx/tx, ny/ty), dim3(tx, ty) >>>
    __global__ void matrixKernel(float* matrix, int width, int height) 
    {
        int x = blockIdx.x * blockDim.x + threadIdx.x;
        int y = blockIdx.y * blockDim.y + threadIdx.y;
        
        if (x < width && y < height) 
        {
            int idx = y * width + x;
            matrix[idx] *= 2.0f;
        }
    }

 

在设计 <<<grid, block>>> 的维度时，需要遵守硬件限制：

参数 硬件上限 实践建议
Block 内最大线程数 1024 个 通常设为 128, 256 或 512
Block 最大维度 (1024, 1024, 64) x * y * z 不超过 1024
Grid 最大维度 (231−1,65535,65535)(2^{31}-1, 65535, 65535) 根据数据量自动计算
Warp 大小 32 个线程 Block 的大小（尤其是一维长度）最好是32 的整数倍，以避免浪费 Warp 内的线程

 

4. 硬件调度

在软件层面，我们将线程组织为 Block。但在 GPU 硬件层面，调度和执行有着自己的规则。

当启动一个 Kernel 时，GPU 会将 Grid 中的 Block 分配给各个 SM 去执行，一个 Block 必须被完整的分配到一个 SM 上，不能跨 SM 分布。

在 SM 内部，Block 并不是作为一个整体同时执行的。SM 会将 Block 中的线程按照 32 个为一组进行划分，每组 32 个线程被称为一个 Wrap：

  • 锁步执行：同一个 Wrap 内的 32 个线程在同一时刻必须执行完全相同的指令
  • 分支发散：如果 Warp 内的线程遇到 if-else 分支，且部分线程走 if，部分走 else，GPU 无法让它们同时执行不同分支。它只能先让走 if 的线程执行（其他线程挂起），然后再让走 else 的线程执行。这会导致性能严重下降，是 CUDA 优化的重点

 

5. 示例：向量加法

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
__global__ void vectorAddKernel(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];
    }
}

void vectorAdd(const float* a, const float* b, float* c, int n)
{
    if (n <= 0)
    {
        return;
    }

    size_t bytes = n * sizeof(float);

    float* dA = nullptr;
    float* dB = nullptr;
    float* dC = nullptr;

    // 1. 在 GPU 设备分配显存
    cudaError_t err = cudaMalloc(&dA, bytes);
    if (err != cudaSuccess)
    {
        throw std::runtime_error(cudaGetErrorString(err));
    }

    err = cudaMalloc(&dB, bytes);
    if (err != cudaSuccess)
    {
        cudaFree(dA);
        throw std::runtime_error(cudaGetErrorString(err));
    }

    err = cudaMalloc(&dC, bytes);
    if (err != cudaSuccess)
    {
        cudaFree(dA);
        cudaFree(dB);
        throw std::runtime_error(cudaGetErrorString(err));
    }

    // 2. 将数据从 Host 拷贝到 Device
    cudaMemcpy(dA, a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(dB, b, bytes, cudaMemcpyHostToDevice);

    // 3. 计算 Block 与 Grid 大小，并启动 Kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vectorAddKernel<<<blocksPerGrid, threadsPerBlock>>>(dA, dB, dC, n);

    // 4. 检查 Launch 错误并同步设备
    err = cudaGetLastError();
    if (err != cudaSuccess)
    {
        cudaFree(dA);
        cudaFree(dB);
        cudaFree(dC);
        throw std::runtime_error(cudaGetErrorString(err));
    }
    cudaDeviceSynchronize();

    // 5. 将计算结果从 Device 拷贝回 Host
    cudaMemcpy(c, dC, bytes, cudaMemcpyDeviceToHost);

    // 6. 释放 Device 显存
    cudaFree(dA);
    cudaFree(dB);
    cudaFree(dC);
}