CUDA核心概念与内存层次

发表于 阅读次数： 本文字数： 5k 阅读时长 ≈ 9 分钟

想在一张图看懂 GPU 的线程组织与内存层次？本文用表格 + Mermaid 图 + 代码示例，3 分钟带你摸清 Kernel、Warp、合并访存与 bank 冲突的底层逻辑。

📚 核心概念速读

名称 一句话定义 省流
Kernel 在 GPU 上并行执行的函数 把循环搬进 GPU
__global__ 修饰 Kernel，CPU 端调用 “这函数在 GPU 上跑”
threadIdx.* / blockIdx.* 线程/块坐标 每线程自带“学号”
<<< grid, block >>> Kernel 启动语法 告诉 GPU “来几拨人干活”
Warp (32 线程) 锁步执行的最小调度单元 小队齐步走 (SIMT)

省流：写 Kernel → 设 Grid/Block → GPU 并行跑 → cudaDeviceSynchronize() 等结果。

1️⃣ Kernel 基础：__global__ 与线程索引

在真正调用 CUDA API 前，建议先写一个通用错误检查宏，否则内核即使失败也可能“静默”通过编译与运行，难以及时发现问题。

1
2
3
4
5
6
7
8
9
10
// 简易错误检查：若调用返回非 cudaSuccess，则打印信息并退出
#define CUDA_CHECK(call)                                                          \
    do {                                                                         \
        cudaError_t err__ = (call);                                              \
        if (err__ != cudaSuccess) {                                              \
            fprintf(stderr, "CUDA Error %s:%d: %s\n", __FILE__, __LINE__,      \
                    cudaGetErrorString(err__));                                  \
            exit(EXIT_FAILURE);                                                  \
        }                                                                        \
    } while (0)

不写 CUDA_CHECK 也能跑，但如果 Launch 参数超界、显存不足或其他运行时错误，就会在 cudaDeviceSynchronize() 处返回非零状态；若从未检查，你会在结果错乱时才发现。宏的成本极低，却能第一时间定位错误，强烈推荐写上。

1
2
3
4
5
6
7
8
9
10
11
__global__ void saxpy(const float* x, const float* y, float* out, float a, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x; // 全局索引
    if (i < n) out[i] = a * x[i] + y[i];          // 每线程做一件小事
}

// host 端调用
int threads = 256;
int blocks  = (n + threads - 1) / threads;
saxpy<<<blocks, threads>>>(d_x, d_y, d_out, 2.0f, n);
CUDA_CHECK(cudaGetLastError());      // 立即检查 launch 错误
CUDA_CHECK(cudaDeviceSynchronize()); // 等待 GPU 跑完
  • __global__：声明 GPU 执行 / CPU 调用 的函数。
  • threadIdx.{x,y,z}：线程在 Block 内 的局部坐标。
  • blockIdx.{x,y,z}：Block 在 Grid 内 的坐标。
  • blockDim.{x,y,z}：Block 尺寸；gridDim.{x,y,z}：Grid 尺寸。

💡 实用技巧：把索引计算抽成 int gid = threadIdx.x + blockIdx.x * blockDim.x;，后续 gid += blockDim.x * gridDim.x 可轻松“网格跨步循环”。

2️⃣ 线程组织：Grid / Block / Thread / Warp



flowchart TD
    G["Grid"] --> B0["Block 0"]
    G --> B1["Block 1"]
    B0 --> T0["Thread 0"]
    B0 --> T1["Thread 1"]
    B0 --> T2["Thread 2"]
    B0 --> T31["Thread 31"]
    subgraph "Warp（32 线程）"
        T0
        T1
        T2
        T31
    end
  • Grid：一次 Kernel Launch 的“全班”；最大可达数十亿线程。
  • Block：可在 SM 间调度迁移的独立单元；内部可用 __syncthreads() 同步。
  • Thread：最细粒度执行体；寄存器私有。
  • Warp (32)：硬件锁步组；SIMT 让 32 线程执行同一指令。

⚠️ 分歧 (Divergence)：同一 Warp 内线程 if/else 走不同分支 → 硬件掩码串行 → 吞吐掉队。

3️⃣ 错误处理与同步

API 作用 建议用法
cudaGetLastError() 捕获 Kernel Launch 级错误 Launch 后立刻调用一次
cudaDeviceSynchronize() CPU 等 GPU 完成 & 捕获运行时错误 调试阶段在关键点加；生产用 Events/Streams 替代

省流：能跑就不代表对——GPU 错误多在 异步 回来，cudaDeviceSynchronize() 是最终裁判。

4️⃣ 内存层次与访存

层次 可见范围 容量 延迟 典型用途
寄存器 线程 KB 级 🟢 最低 私有临时变量
Shared Block 64~100 KB 🟢 低 Block 内数据复用
L2 Cache 全 GPU MB 级 🟡 中 缓解全局访存
Global 全 GPU GB 级 🔴 高 大规模数据
Constant / Texture 全 GPU 64 KB / 48 KB 🟢(只读缓存) 只读 & 空间局部性 


flowchart TD
    Reg[寄存器] --> Shm[共享内存]
    Shm --> L2[L2 Cache]
    L2 --> Mem[全局显存]
    Mem --> Host[主机内存]

为什么表格里有 Constant / Texture，却没画在图上？
这两块只读缓存（常量缓存 & 纹理缓存）实际上挂在 L2 旁边，走的是同一显存通道。为了让主数据路径更直观，我在图里省略了它们；你可以把它想成 “贴在 L2 侧面的专用快速通道”，专门服务 小而频繁的只读数据（比如卷积核、查找表）。

5️⃣ 合并访存 (Coalescing)

同一个 Warp 的 32 线程若 按地址连续 访问，可被 GPU 合并为 1~2 次 32/64/128B 事务：

1
2
3
4
// 好：线程 i 访问 a[i]
float v = a[threadIdx.x + blockIdx.x * blockDim.x];

// 坏：线程 i 访问 a[i * stride]

省流：相邻线程 → 相邻地址 → 带宽“打包发货”🚚。

6️⃣ 共享内存 Bank 冲突与 +1 Padding

共享内存就像 32 条并排的高速收费通道（32 个 bank）。每条通道一次只能通过 1 辆车（一次访问）。

📦 tile 是什么？为什么和 shared 形影不离？

  • tile 的定义：把大矩阵/张量切成 TILE_DIM × TILE_DIM 的“小方块”，同一个 thread block 负责读取、处理并写回这一方块。这个子区域就叫 tile。
  • 与 shared 的关系：tile 通常先被搬进 shared memory 暂存，线程在片上对其做转置/卷积/归约等操作，再合并写回全局内存。
  • 为何配合使用
    1. 合并访存：按列的零散全局访问 → 读入 tile → 在 shared 内转置 → 按行合并写，带宽翻倍。
    2. 数据复用：tile 内数据可被 block 中不同线程反复用（矩阵乘、卷积滑窗）。
    3. 避开 Bank 冲突：tile 往往声明为 tile[T][T+1]，多出的 +1 用来打散 32-bank 映射。
  • 典型模板： 
    1
    2
    3
    4
    5
    6
    7
    8
    9
    constexpr int TILE_DIM = 32, BLOCK_ROWS = 8;
    __global__ void kernel(const float* A, float* B, int M, int N) {
        __shared__ float tile[TILE_DIM][TILE_DIM + 1]; // tile 存在 shared
        int x = blockIdx.x * TILE_DIM + threadIdx.x;
        int y = blockIdx.y * TILE_DIM + threadIdx.y;
        // 1) 全局 → tile (合并读)
        // 2) __syncthreads(); 在片上处理/转置
        // 3) tile → 全局 (合并写)
    }
  • 选型注意：TILE_DIM=32, BLOCK_ROWS=8 是常用配置；tile 过大将占满 shared 或寄存器，影响 Occupancy，需要权衡。

TILE_DIM 通常对应 blockDim.x（横向线程数 & tile 宽度）。 BLOCK_ROWS 对应 blockDim.y，代表一次搬运/计算的 行数批次；典型 8 行，通过循环搬满 32 行。

  • 理想：不同线程去不同通道 → 并行一次性通过 🟢。
  • 冲突：多辆车挤同一通道 → 排队串行，性能打折 🔴。

常见触发 Bank 冲突的模式（未做 +1 时）
1. 同列访问：矩阵转置/列拷贝，线程 (i,j 固定) → 地址步长 = 32×4B。
2. Stride 为 32 的倍数：Tile 宽度 32、64… 行间起点对齐同 bank。
3. 跨步写入：线程 i 访问 base + i*32 等模 32 重合的地址。
4. 纯行或纯列归约：只动 x 或 y，导致多线程集中一个 bank。
5. 稀疏 gather/scatter：索引映射使 (addr/4)&31 分布极度不均。

+1 Padding 的效果？
• ✅ 有效：前 3 类（同列访问、stride 为 32 的倍数、跨步 32）——本质都是 行间步长=32k，加 1 立刻破坏同 bank 对齐。
• ❌ 无效：纯行/列归约与高随机 gather/scatter——冲突来源是线程集中到同一地址或分布不可控，单靠填充无法分散，需要换算法或分段访问。

1
__shared__ float tile[32][33]; // 33 = 32 + 1   (列方向额外 +1 填充)

为什么 +1 就能避开冲突？

  1. tile[行][列] 在显存里按 行主序 排布。
  2. 第 0 行开始地址对齐 bank0，第 1 行本该也落 bank0。
  3. 把列宽从 32 ➡️ 33，让下一行起始地址向后多挪 4B，落到 bank1。
  4. 结果：32 行分别落到 32 个 bank，所有线程同时访问首元素 → 零冲突。

Bank 冲突为什么只看 Warp 内同周期？ GPU 每个时钟只让 一个 Warp 的共享内存指令发射 → 冲突检测仅在这 32 线程之间进行；下一 Warp 要等上一 Warp 完成后才发射，时间上已错开，不会相互叠加。

7️⃣ 主机 ↔ 设备拷贝

拷贝链路可以类比 “快递寄包裹”：

货物打包方式 举例 “物流专用车道” 带宽 可边走边干活？
可分页内存 (pageable) 普通 new/malloc 普通公路 中速 否，需要等车到站
固定页内存 (pinned) cudaHostAlloc() / cudaMallocHost() 高速专用线 高速 否，但上车更快
Pinned + Async Stream cudaMemcpyAsync 高速专用线 + 双车道 高速 ✅ 运输 & 计算并行
  • 为什么 pinned 更快？ 省去 OS 页锁定 & 临时缓冲，DMA 可直接搬运。
  • 为什么要异步？ 像“发快递同时在家里干别的”——GPU 计算与 PCIe 传输并行，流水线利用率拉满。
  • 可分页内存适合谁？ 小脚本、偶尔拷 100 KB 的日志。
  • Pinned but 同步拷贝：数据量大但计算无法并行，至少省掉锁页开销。
  • Pinned + Async Stream：训练深度学习模型、实时视频批处理这类“边算边吃数据”的场景首选。

🔎 调试小抄

  1. cuda-memcheck ./a.out：越界/初始化等内存错误。
  2. nvprof --print-gpu-trace ./a.out：粗粒度 Timeline。
  3. Nsight Compute / Nsight Systems：瓶颈定位。

📝 小结

  • Kernel 三件套：__global__ + 线程索引 + <<<grid, block>>>。
  • 线程组织公式：Warp(32) ⊂ Block ⊂ Grid，分歧少、Occupancy 高。
  • 访存第一性原理：相邻访问 + Tile 复用 + 避免 bank 冲突。
  • 拷贝优化：Pinned + Async Stream，让 PCIe 传输“打多路复用”。

记住一句话：GPU 用“成群线程”换吞吐，用“切 Warp”藏延迟，性能归根结底是访存模式与并行度的艺术 ✨