gpu-memory-model

GPU 内存模型

目的

说明供优化内核的 Agent 使用的 GPU 执行与内存模型：SIMT 执行、warp（32）与 wavefront（64）的分歧开销、全局内存合并规则、共享内存 bank 冲突、L1/L2 缓存行为、原子内存排序，以及占用率与延迟隐藏之间的权衡。

适用场景

• 诊断理论带宽很高但内核仍受内存限制的原因
• 了解分支导致的 warp 分歧
• 修复分块算法中的共享内存 bank 冲突
• 选择 block 大小，以权衡占用率与寄存器压力
• 在 NVIDIA（warp 32）与 AMD（wavefront 64）之间移植内核
• 分析并行归约中的原子操作竞争

工作流

1. SIMT 执行模型

GPU hardware
├── Device
│   └── SM / CU (Streaming Multiprocessor / Compute Unit)
│       ├── Warp schedulers (NVIDIA) or Wavefront schedulers (AMD)
│       │   └── Warp/Wavefront (32 or 64 threads in lockstep)
│       ├── Register file (partitioned per thread)
│       ├── Shared memory / LDS (per SM)
│       └── L1 cache (often shared with shared memory)
└── L2 cache (device-wide) → DRAM/HBM

SIMT（单指令多线程）：一条指令流驱动一个 warp/wavefront；每个线程都有自己的寄存器和线程 ID，但会锁步执行相同的指令。

2. Warp 与 wavefront

厂商	单元大小	名称
NVIDIA	32 个线程	Warp
AMD	64 个线程	Wavefront

影响：

• 归约树：NVIDIA 从 16→8→4→2→1 逐步减半；AMD 则从 32→16→8→4→2→1
• Block 大小：优先选择 32（NVIDIA）或 64（AMD）的倍数
• 占用率计数器会报告每个 SM 上活跃的 warp/wavefront 数量

3. Warp 分歧开销模型

当一个 warp 中的线程进入不同分支时，硬件会串行执行各个路径：

// Divergent: half warp does A, half does B → 2x instruction issue
if (threadIdx.x % 2 == 0) {
    result = expensive_a(data[idx]);
} else {
    result = expensive_b(data[idx]);
}

// Non-divergent: all threads same path
result = expensive_a(data[idx]);

缓解方法：

• 谓词执行：同时计算两种结果，再用 ?: 选择（以计算量换取执行一致性）
• 为不同代码路径使用独立内核
• 仅基于 block 级数据进行分支（在 warp 内保持一致）
• 遍历各个 bin，而不是让每个线程都根据 if (data[i] < threshold) 分支，并产生分散的结果

分歧开销 ≈ 所有已执行路径的开销之和（而非最大值）。

4. 全局内存合并

NVIDIA 的合并规则（简化版）：一个 warp 中的线程访问连续的 4 字节字时，会产生单个 128 字节事务。

// Coalesced: consecutive threads → consecutive addresses
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];

// Uncoalesced: stride access
float val = data[threadIdx.x * stride];  // stride > 1

// Partially coalesced: misaligned start
float val = data[base + threadIdx.x * 3];

AoS 与 SoA 的影响：

// AoS — poor coalescing when reading one field
struct Particle { float x, y, z; };
float x = particles[i].x;  // threads read with stride 3

// SoA — coalesced
float x = pos_x[i];

5. 共享内存 bank 冲突

共享内存分为 32 个 bank（每个包含 4 字节字）。同时访问同一 bank 中不同地址时，会发生串行化。

__shared__ float tile[32][32];

// Bank conflict: all threads access tile[threadIdx.x][0]
// 32 threads, 32 banks, but column 0 → same bank per row offset
float val = tile[threadIdx.x][0];

// Fix: pad columns to break bank alignment
__shared__ float tile[32][33];  // +1 padding

检测方法：使用 Nsight Compute 指标 l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum 或 NCU 共享内存加载冲突指标。

6. L1/L2 缓存行为

级别	范围	说明
L1	每个 SM	通常与共享内存统一配置；可调整分配比例
L2	整个设备	缓存行通常为 128 字节
纹理/L1 只读缓存	每个 SM	为访问一致的数据提供缓存只读路径

有利于缓存的模式：

• 空间局部性：相邻线程访问相邻内存
• 时间局部性：在重新从全局内存读取前，先在共享内存中复用数据
• 避免随机分散访问：原子更新和指针追逐会降低缓存效率

// Cache-friendly tile load
for (int t = 0; t < num_tiles; t++) {
    __shared__ float smem[TILE][TILE];
    smem[ty][tx] = global[row * N + t * TILE + tx];
    __syncthreads();
    // compute from smem — L1/L2 only hit on first load per tile
}

7. 原子操作与内存排序

GPU 原子操作（atomicAdd、atomicCAS、atomicExch）可确保以同一地址为目标的线程之间具有顺序一致性，但高竞争会导致执行串行化。

// Bad: all threads atomic to one counter
atomicAdd(&global_sum, local_val);

// Better: per-block reduction, one atomic per block
__shared__ float block_sum;
// ... warp reduce to block_sum ...
if (threadIdx.x == 0)
    atomicAdd(&global_sum, block_sum);

HIP/CUDA 内存栅栏：

__threadfence_block();  // visible to threads in same block
__threadfence();        // visible to all threads on device
__threadfence_system(); // visible to host (expensive)

8. 占用率与延迟隐藏

Occupancy tradeoff
├── High occupancy → more warps to hide memory latency
│   └── Costs: fewer registers/SM, less shared mem per block
└── Low occupancy + high ILP → enough independent instructions per warp
    └── Works for compute-bound kernels with deep pipelines

决策树：

Memory-bound kernel?
├── Yes → maximize active warps (occupancy), coalesce, tile with shared mem
└── No (compute-bound) → may lower occupancy if registers enable more ILP

# Measure achieved occupancy
ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active ./app

经验法则：受内存限制的内核需要至少 50% 的占用率；如果指令级并行度充足，计算受限的内核在 25% 的占用率下也可能运行良好。

9. 快速参考表

症状	可能原因	检查项
DRAM 吞吐量低	访问未合并	NCU 内存工作负载分析
共享内存加载停顿	Bank 冲突	对数组进行填充；更改访问模式
barrier 停顿高	缺少 __syncthreads 或 barrier 处发生分歧	synccheck
原子操作瓶颈	竞争激烈的原子操作过多	分层归约
占用率低	寄存器/共享内存压力	launch__occupancy_limit_* 指标

常见问题

症状	原因	修复方法
SoA→AoS 更改后速度慢了 10 倍	跨步合并访问被破坏	将热点字段保留在 SoA 布局中
分块矩阵乘法比朴素实现更慢	共享内存分块中存在存储体冲突	对共享数组的列进行填充
代码相同，但 NVIDIA 和 AMD 上的性能不同	Warp 为 32，而 wavefront 为 64	重新调整块大小和归约操作
原子操作结果正确，但速度较慢	全局争用	先进行块级归约
L2 命中率高，但速度仍然较慢	L2 带宽已饱和	减少传输的总字节数

相关技能

• skills/gpu/cuda — 使用此模型的实用内核模式
• skills/gpu/cuda-profiling — 用于验证合并访问和占用率的指标
• skills/gpu/hip-rocm — AMD wavefront 专属调优
• skills/low-level-programming/cpu-cache-opt — CPU 缓存概念（类似）
• skills/low-level-programming/simd-intrinsics — CPU 端的向量宽度
• skills/profilers/hardware-counters — 通用缓存未命中测量