枕石的个人博客
文章 标签 分类 关于
枕石的个人博客

Contents

CUDA编程基础

 included in CUDA
   4938 words   10 minutes 

CUDA（Compute Unified Device Architecture）是 NVIDIA 推出的并行计算平台和编程模型。它允许程序员使用类似 C/C++ 的语法编写在 GPU 上运行的函数，将同一项计算分配给大量线程并行完成。本文不会只罗列 CUDA API，而是沿着“GPU 硬件 → CUDA 线程模型 → Kernel 调度 → 内存层次 → 基础优化”这条主线，建立一个完整的 CUDA 编程认知。

一、从 CPU 到 GPU：为什么需要 CUDA

CPU 和 GPU 都能执行计算，但两者的设计目标不同。

CPU 使用少量强大的核心，依靠复杂的控制逻辑、大容量缓存和分支预测等机制，尽可能降低单个任务的延迟。GPU 则把更多芯片面积用于计算单元，通过大量并行线程获得更高的吞吐量。

特性CPUGPU
主要目标低延迟、复杂控制高吞吐、数据并行
核心特点数量少，单核能力强计算单元多，适合批量计算
适合任务分支复杂、串行依赖强计算规则、数据量大
常见场景操作系统、业务逻辑图形渲染、科学计算、深度学习

例如，对两个包含百万个元素的数组做加法，每个结果元素之间没有依赖：

C[0] = A[0] + B[0]
C[1] = A[1] + B[1]
C[2] = A[2] + B[2]
...

这类任务很适合拆分成大量线程并行执行。CUDA 的作用，就是向程序员提供一套描述这种并行任务的软件模型。

GPU 的整体硬件结构

NVIDIA GPU 的具体结构会随架构代际而变化，但可以用下面的层次建立基本认识：

GPU
├── 前端与任务分发单元
├── GPC（Graphics Processing Cluster）
│   └── TPC（Texture Processing Cluster）
│       └── SM（Streaming Multiprocessor）
├── 片上互连与 L2 Cache
├── Memory Controller
└── HBM / GDDR 显存

SM（Streaming Multiprocessor，流式多处理器）是执行 CUDA 程序的主要硬件单元。一个 SM 不只包含 CUDA Core，还包含 Warp Scheduler、Register File、Load/Store Unit、SFU、Tensor Core 以及 L1 Cache / Shared Memory 等资源。

┌────────────────── SM ──────────────────┐
│ Warp Scheduler / Dispatch Unit         │
│ Register File                         │
│ FP32 / INT32 / FP64 执行流水线          │
│ Tensor Core / LD-ST Unit / SFU        │
│ L1 Data Cache / Shared Memory         │
└─────────────────────────────────────────┘

/images/CUDA编程/image-3.png

NVIDIA GPU 整体架构图， GPC、TPC、SM、L2 Cache、Memory Controller 和 HBM。

/images/CUDA编程/image-4.png

SM 内部结构图，Warp Scheduler、Register File、各类执行单元以及 L1 / Shared Memory。

需要特别注意：SM 不等于 CUDA Core，CUDA Thread 也不会永久绑定到某一个 CUDA Core。线程会先组成 Warp，再由 SM 中的调度器将指令发射给对应的执行单元。

二、CUDA 编程模型：Kernel、Grid、Block 和 Thread

CUDA 采用异构编程模型：

  • Host：CPU 及 CPU 内存；
  • Device：GPU 及 GPU 显存；
  • Kernel：由 Host 发起、在 Device 上并行执行的函数。

一个典型 CUDA 程序的执行流程如下：

CPU 分配和准备数据
        ↓
分配 GPU 显存
        ↓
将数据从 Host 复制到 Device
        ↓
启动 Kernel
        ↓
GPU 并行执行 Kernel
        ↓
将结果从 Device 复制回 Host
        ↓
释放资源

Kernel

Kernel 使用 __global__ 修饰，调用时需要在函数名和参数之间指定执行配置：

__global__ void kernel_name(/* parameters */)
{
    // 每个 CUDA Thread 都会执行这段代码
}

kernel_name<<<num_blocks, threads_per_block>>>(/* arguments */);

<<<num_blocks, threads_per_block>>> 表示这次 Kernel 启动使用多少个 Block，以及每个 Block 中有多少个 Thread。

Grid、Block 和 Thread

CUDA 将一次 Kernel 启动创建的并行任务组织成以下层次：

Kernel Launch
    └── Grid
        ├── Block 0
        │   ├── Thread 0
        │   ├── Thread 1
        │   └── ...
        ├── Block 1
        └── ...
  • Thread 是基本的逻辑并行单位。每个 Thread 都会从头到尾执行一次 Kernel，但可以根据自己的索引处理不同数据。
  • Block 是一组 Thread，是资源驻留和线程协作的重要边界。同一 Block 的线程位于同一 SM，可以共享 Shared Memory，并通过 __syncthreads() 同步。
  • Grid 是一次 Kernel 启动产生的所有 Block 的集合。Grid 的 Block 数量可以远大于 GPU 的 SM 数量。

CUDA 为每个线程提供了几个内置变量：

内置变量含义
threadIdxThread 在当前 Block 中的索引
blockIdxBlock 在当前 Grid 中的索引
blockDim每个 Block 的维度
gridDimGrid 的 Block 维度

在一维数据中，线程的全局索引通常计算为：

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

例如，blockDim.x = 256、blockIdx.x = 3、threadIdx.x = 5 时：

i = 3 × 256 + 5 = 773

该线程就可以负责计算数组的第 773 个元素。Grid 和 Block 还可以是二维或三维结构，便于表示图像、矩阵和三维网格等数据。

/images/CUDA编程/image-5.png

CUDA Grid、Block 和 Thread 的三维层次图。

三、第一个 CUDA 程序：向量加法

下面用向量加法串起 CUDA 程序的基本步骤。为了让主要流程更清楚，示例使用一个宏统一检查 CUDA Runtime API 的返回值。

#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>

#define CUDA_CHECK(call)                                             \
    do {                                                             \
        cudaError_t error = (call);                                  \
        if (error != cudaSuccess) {                                  \
            fprintf(stderr, "CUDA error at %s:%d: %s\n",           \
                    __FILE__, __LINE__, cudaGetErrorString(error));  \
            exit(EXIT_FAILURE);                                      \
        }                                                            \
    } while (0)

__global__ void vector_add(const float *a,
                           const 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(void)
{
    const int n = 1 << 20;
    const size_t bytes = n * sizeof(float);

    float *h_a = (float *)malloc(bytes);
    float *h_b = (float *)malloc(bytes);
    float *h_c = (float *)malloc(bytes);

    if (h_a == NULL || h_b == NULL || h_c == NULL) {
        fprintf(stderr, "Host memory allocation failed\n");
        free(h_a);
        free(h_b);
        free(h_c);
        return EXIT_FAILURE;
    }

    for (int i = 0; i < n; ++i) {
        h_a[i] = (float)i;
        h_b[i] = (float)(2 * i);
    }

    float *d_a = NULL;
    float *d_b = NULL;
    float *d_c = NULL;

    CUDA_CHECK(cudaMalloc((void **)&d_a, bytes));
    CUDA_CHECK(cudaMalloc((void **)&d_b, bytes));
    CUDA_CHECK(cudaMalloc((void **)&d_c, bytes));

    CUDA_CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));

    const int threads_per_block = 256;
    const int num_blocks = (n + threads_per_block - 1)
                         / threads_per_block;

    vector_add<<<num_blocks, threads_per_block>>>(d_a, d_b, d_c, n);

    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());

    CUDA_CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));

    printf("c[100] = %.0f\n", h_c[100]);

    CUDA_CHECK(cudaFree(d_a));
    CUDA_CHECK(cudaFree(d_b));
    CUDA_CHECK(cudaFree(d_c));
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

使用 nvcc 编译并运行：

nvcc -O2 -o vector_add vector_add.cu
./vector_add

输出：

c[100] = 300

计算 Block 数量

向上取整公式：

int num_blocks = (n + threads_per_block - 1) / threads_per_block;

它保证了即使 n 不能被 threads_per_block 整除，也会创建足够的线程覆盖全部元素。因为最后一个 Block 中可能包含多余线程，Kernel 内必须进行边界检查：

if (i < n) {
    c[i] = a[i] + b[i];
}

Kernel 启动与错误检查

Kernel 启动相对 Host 线程通常是异步的，即 CPU 提交 Kernel 后可以继续执行。因此，下面两步检查的问题不完全相同：

CUDA_CHECK(cudaGetLastError());       // 检查 Kernel 启动配置等错误
CUDA_CHECK(cudaDeviceSynchronize());  // 等待执行完成，暴露异步执行错误

在真实程序中，也可以使用 Stream 和 Event 表达更精细的依赖，而不是每次都调用全设备同步。

四、从 Block 到 Warp：Kernel 如何在 SM 上执行

CUDA API 直接暴露的是 Grid、Block 和 Thread，但 GPU 实际调度时还有一个关键层次：Warp。

Grid
  ↓
Block
  ↓
Warp
  ↓
Thread

Block 如何分配到 SM

当 CPU 启动 Kernel 后，GPU 会将 Grid 中的 Block 不断分配给有可用资源的 SM。假设一个 GPU 有 4 个 SM，Grid 中有 10 个 Block，某一时刻可能是：

SM 0: Block 0, Block 4
SM 1: Block 1, Block 5
SM 2: Block 2, Block 6
SM 3: Block 3, Block 7

Waiting: Block 8, Block 9

当某个 Block 执行完成并释放资源后，等待中的 Block 再进入该 SM。由此可以得到三个重要结论：

  1. 一个 Block 的所有线程会驻留在同一个 SM 上；
  2. 一个 SM 在资源足够时可以同时驻留多个 Block；
  3. Block 与 SM 没有固定的编号对应，不同 Block 的执行顺序也不应被程序依赖。

一个 SM 能同时容纳多少 Block，主要受以下资源限制：

  • 每个 Block 的 Thread 数；
  • 每个 Thread 使用的寄存器数；
  • 每个 Block 使用的 Shared Memory；
  • 架构支持的最大驻留 Block 和 Warp 数。

寄存器或 Shared Memory 用量过大，都可能减少同时驻留的 Block 和 Warp 数量。

Warp 与 SIMT

Block 进入 SM 后，会被划分为 Warp。在 CUDA 中：

1 Warp = 32 Threads

例如，一个 Block 包含 256 个 Thread：

Warp 0: Thread   0 ~  31
Warp 1: Thread  32 ~  63
Warp 2: Thread  64 ~  95
...
Warp 7: Thread 224 ~ 255

GPU 以 SIMT（Single Instruction, Multiple Threads）方式执行 Warp。程序员仍然为每个 Thread 编写独立的程序逻辑，每个 Thread 也有自己的索引、寄存器状态和控制流；但在硬件上，同一 Warp 的线程被组织在一起执行指令。

Warp Scheduler 与延迟隐藏

SM 通常会同时保持多个已驻留 Warp。当某个 Warp 等待显存数据或指令依赖时，Warp Scheduler 可以选择其他已就绪的 Warp 执行：

Warp 0: 等待 Global Memory
Warp 1: Ready  → 执行 FP32 指令
Warp 2: 等待数据依赖
Warp 3: Ready  → 执行 Load/Store 指令

这种机制并没有消除访存延迟，而是用其他 Warp 的有效工作覆盖等待时间，通常称为延迟隐藏（Latency Hiding）。

Occupancy 表示 SM 上活跃 Warp 数与架构允许的最大活跃 Warp 数之比。较多的活跃 Warp 往往有助于隐藏延迟，但** Occupancy 越高并不代表性能一定越高**。为追求 Occupancy 而过度减少寄存器，可能导致额外访存，反而降低性能。

Warp Divergence

当同一 Warp 内的线程走向不同分支时，会产生 Warp Divergence：

if (threadIdx.x % 2 == 0) {
    y = function_a();
} else {
    y = function_b();
}

概念上，执行 if 分支时，不走该路径的线程会被屏蔽；然后再执行 else 分支。两条分支都需要执行，有效计算单元利用率因此下降。分歧并不会导致计算错误，但在热点路径上应尽量让同一 Warp 中的线程执行相似的控制流。

五、CUDA 内存模型

GPU 性能不仅取决于计算单元数量，也取决于数据能否高效地送到计算单元。CUDA 程序员需要区分不同存储层次的作用域、管理方式和性能特征。

存储资源典型作用域位置管理方式主要特点
RegisterThreadSM 片上编译器延迟低，数量有限
Local MemoryThreadDevice Memory编译器线程私有，但不是片上高速内存
Shared MemoryBlockSM 片上程序员Block 内共享，延迟低
L1 CacheSMSM 片上主要由硬件管理每个 SM 私有
L2 CacheGPUGPU 片上主要由硬件管理所有 SM 共享
Global MemoryDeviceHBM / GDDR程序员分配容量大，延迟高
Constant MemoryGrid 只读Device Memory，有专用缓存程序员适合线程读取相同的少量常量

Register 和 Local Memory

Kernel 中的局部标量变量通常会被编译器放入寄存器：

int i = blockIdx.x * blockDim.x + threadIdx.x;
float value = input[i] * 2.0f;

i 和 value 可能使用寄存器，但 C/C++ 局部变量并不保证一定存入寄存器。大型局部数组、无法在编译时确定索引的数组，或寄存器不足时的 spill，都可能使用 Local Memory。

local 表示它在逻辑上属于单个 Thread，不表示它位于 SM 内部的高速存储器。Local Memory 实际位于 Device Memory 空间，只是可以通过缓存改善部分访问。

Shared Memory

Shared Memory 是由程序员显式管理的 Block 内共享存储空间：

__global__ void reverse(const float *input, float *output)
{
    __shared__ float buffer[256];

    int local_id = threadIdx.x;
    int global_id = blockIdx.x * blockDim.x + local_id;

    buffer[local_id] = input[global_id];
    __syncthreads();

    output[global_id] = buffer[blockDim.x - 1 - local_id];
}

buffer 属于整个 Block，同一 Block 中的所有 Thread 都可访问。__syncthreads() 保证 Block 内的线程都完成写入后，再开始下一阶段读取。

Shared Memory 和 L1 Cache 的底层存储资源在某些架构上是统一的，但它们的编程语义不同：

  • Shared Memory 的数据放置、读写和同步由程序员控制；
  • L1 Cache 是 Global / Local Memory 的缓存层次，主要由硬件管理。

不同 Block 之间不能通过普通 Shared Memory 共享数据，也不能将 __syncthreads() 当作 Grid 级同步。

Global Memory 与缓存

Global Memory 通常对应 GPU 的 HBM 或 GDDR 显存。它容量大，所有 SM 都可以访问，但相比寄存器和 Shared Memory 有更高延迟。典型数据路径可以简化为：

Register / Shared Memory
          ↑
       L1 Cache
          ↑
       L2 Cache
          ↑
   HBM / GDDR 显存

CUDA 中没有名为“SRAM”的可直接申请编程空间。更准确的说法是：程序员可显式使用 Shared Memory 和 Global Memory，寄存器由编译器分配，L1/L2 Cache 主要由硬件管理；SRAM 和 DRAM 是更接近硬件实现的存储介质分类。

图片待补： CUDA 内存层次图，标出 Thread 私有的 Register / Local Memory、Block 共享的 Shared Memory、每个 SM 的 L1 Cache、GPU 共享的 L2 Cache 和 Global Memory。

六、基础性能优化与常见误区

掌握线程和内存模型后，CUDA 优化的主要问题可以归结为：让计算单元尽量忙碌，同时减少慢速、低效的数据搬运。

Global Memory 合并访问

当一个 Warp 中的相邻线程访问连续内存地址时：

Thread 0  → data[0]
Thread 1  → data[1]
Thread 2  → data[2]
...
Thread 31 → data[31]

硬件能够用较少的 Memory Transaction 完成访问，这称为合并访问（Coalesced Access）。如果相邻线程访问大跨度、高度分散的地址，就可能需要更多交易，降低有效带宽。

Shared Memory Tiling

对会反复使用的数据，可以让 Block 内的线程协作将它从 Global Memory 搬到 Shared Memory，然后在片上重复使用：

HBM / Global Memory
        ↓  协作加载
Shared Memory Tile
        ↓  重复使用
Registers
        ↓
ALU / Tensor Core

矩阵乘法、Reduction、Softmax 和 Attention 等高性能 Kernel 都大量使用这个思路。

Shared Memory Bank Conflict

Shared Memory 由多个 Bank 组成。当同一次访问中的多个线程以不利方式访问同一 Bank 时，访问可能被拆分并串行化，这称为 Bank Conflict。

经典例子是矩阵转置。对于：

__shared__ float tile[32][32];

某些按列访问模式会产生严重冲突。通常可以增加一列 Padding 改变地址到 Bank 的映射：

__shared__ float tile[32][33];

常见误区

  1. 一个 Thread 对应一个 CUDA Core：错误。Thread 先组成 Warp，再由调度器将指令发射给对应的执行流水线，不存在永久一对一绑定。
  2. 一个 Block 对应一个 SM：不完整。一个 Block 只会驻留在一个 SM 上，但一个 SM 可能同时驻留多个 Block。
  3. Local Memory 是高速片上内存：错误。local 表示 Thread 私有作用域，它实际位于 Device Memory 空间。
  4. Shared Memory 可以被不同 Block 共享：对普通 CUDA Block 模型不成立。Shared Memory 的常规作用域是 Block。
  5. Kernel 启动后 CPU 会立即等待：错误。Kernel 启动相对 Host 线程通常异步，需要通过同步调用、Stream 或 Event 管理依赖。
  6. Occupancy 越高性能一定越好：错误。性能还受访存模式、计算强度、缓存命中率、指令依赖、Warp Divergence 和 Bank Conflict 等因素影响。

小结

学习 CUDA 的关键不是背诵几个 Runtime API，而是理解下面这条完整路径：

Kernel Launch
    → Grid
    → Block 分配到 SM
    → Block 划分为 Warp
    → Warp Scheduler 发射指令
    → 执行单元计算
    → Register / Shared Memory / Global Memory 之间搬运数据

建立这个模型后，再学习 Reduction、GEMM、Softmax、FlashAttention 和 PagedAttention 等算法时，就能够从线程组织、数据复用和存储层次的角度理解其优化动机。

参考资料

  1. CUDA Programming Guide
  2. CUDA C++ Best Practices Guide
  3. NVIDIA A100 Tensor Core GPU Architecture
  4. CUDA Refresher: The CUDA Programming Model
Updated on 2026-05-24
 CUDA
Back | Home