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>>> 语法来定义 gridDimblockDim

线程组织的核心目的,就是计算出当前线程在整个 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;
    }
    }