01讲了谁在执行线程层次这一课讲数据放在哪内存层次。GPU kernel 性能的 80% 取决于你如何管理内存而不是计算本身。因为 GPU 的计算单元极快但内存访问相对很慢——如何喂数据给计算单元就是优化的核心。GPU 内存全景图速度 ↑ 容量 ↓ ┌─────────────────────────────────────────────────────────┐ │ Register寄存器 每线程私有 ~1 cycle 延迟 │ ← 最快 │ 数量每 SM 65536 个 32-bit │ ├─────────────────────────────────────────────────────────┤ │ Shared Memory 每 Block 共享 ~20-30 cycles │ ← 用户可控的 L1 │ 共享内存 大小48-228 KB / SM │ ├─────────────────────────────────────────────────────────┤ │ L1 Cache 每 SM 私有 ~30-80 cycles │ ← 硬件自动管理 │ 与 Shared Memory 共享物理 SRAM │ ├─────────────────────────────────────────────────────────┤ │ L2 Cache 全 GPU 共享 ~200-800 cycles │ │ A100: 40 MB, H100: 50 MB │ ├─────────────────────────────────────────────────────────┤ │ Global Memory全局显存 所有 SM 共享 ~300-600 cycles │ ← 最慢 │ A100: 80 GB HBM2e │ │ 带宽: ~2 TB/s │ └─────────────────────────────────────────────────────────┘ 速度 ↓ 容量 ↑cycle的全称是 Clock Cycle时钟周期在计算机体系结构和 GPU 编程中它是衡量时间/延迟的最基本单位。你可以把它理解为 GPU 内部节拍器打出的 “一个节拍”。速度差距Global Memory 的一次访问够 Register 执行几百条指令。这就是为什么你需要把频繁使用的数据搬到 Shared Memory 或 Register 里。各层内存的属性总结内存类型作用域生命周期速度容量谁管理Register单个线程kernel 执行期间~1 cycle255 个/线程编译器Shared Memory一个 Block 内所有线程Block 执行期间~20 cycles48-228 KB/SM程序员手动L1 Cache单个 SM硬件决定~30 cycles与 SM 共享 SRAM硬件自动L2 Cache全 GPU硬件决定~200 cycles40-50 MB硬件自动Global Memory全 GPU程序员分配~300 cycles数十 GB程序员分配Global Memory 与 Coalesced Access合并访问Global Memory 是最慢但容量最大的内存几乎所有数据最初都在这里。如何访问它直接决定了 kernel 的带宽利用率。什么是 Coalesced AccessGPU 的 Global Memory 读写不是按单个线程发起的而是以warp32 个线程为单位发起的。一个 warp 访问 Global Memory 时硬件会把 32 个线程的地址合并成尽量少的内存事务memory transaction。完美合并Coalesced32 个线程访问连续内存地址硬件只需1 次128 字节的内存事务。Thread 0 → addr[0] ┐ Thread 1 → addr[1] │ Thread 2 → addr[2] ├─→ 合并为 1 次 128B 事务 ✓ ... │ Thread 31 → addr[31] ┘非合并Uncoalesced32 个线程访问分散地址硬件需要32 次独立内存事务。Thread 0 → addr[0] ┐ Thread 1 → addr[1024] │ Thread 2 → addr[2048] ├─→ 32 次独立事务 ✗ 带宽浪费 ~30x ... │ Thread 31 → addr[31744] ┘关键规则对于一维数组threadIdx.x相邻的线程访问相邻内存地址 → 完美合并。// ✓ 完美合并连续线程访问连续地址 int tid threadIdx.x blockIdx.x * blockDim.x; float val array[tid]; // thread 0→[0], thread 1→[1], ... // ✗ 非合并步长访问 int tid threadIdx.x blockIdx.x * blockDim.x; float val array[tid * 32]; // thread 0→[0], thread 1→[32], thread 2→[64]...二维数组的访问陷阱矩阵运算必知C/CUDA 中矩阵是行优先存储row-majormatrix[row][col]在内存中是matrix[row * num_cols col]。// ✓ 合并访问threadIdx.x 对应列同一行内连续 int row threadIdx.y blockIdx.y * blockDim.y; int col threadIdx.x blockIdx.x * blockDim.x; float val matrix[row * num_cols col]; // 同一 warp 内 threadIdx.x 连续 → col 连续 → 地址连续 → 合并 ✓ // ✗ 非合并threadIdx.x 对应行跨行访问 int row threadIdx.x blockIdx.x * blockDim.x; int col threadIdx.y blockIdx.y * blockDim.y; float val matrix[row * num_cols col]; // 同一 warp 内 threadIdx.x 连续 → row 连续 → 地址跳跃 num_cols → 不合并 ✗结论操作矩阵时让threadIdx.x对应列方向内存连续方向。这在写 matmul、attention 等 kernel 时至关重要。对齐要求起始地址最好是128 字节对齐即 32 个 float 的整数倍否则一次 warp 访问可能跨越两个 128B segment需要 2 次事务。Shared Memory共享内存Shared Memory 是 CUDA 优化最重要的武器。它位于 SM 芯片上速度接近寄存器容量有限但远快于 Global Memory。为什么需要 Shared Memory核心场景Block 内多个线程需要读取相同的数据。如果每个线程都从 Global Memory 读就浪费了大量带宽。把数据先搬到 Shared Memory然后所有线程从 Shared Memory 读——带宽节省 N 倍N Block 大小。不使用 Shared Memory: Global Memory → Thread 0 读 data[i] Global Memory → Thread 1 读 data[i] ← 重复读取浪费带宽 Global Memory → Thread 2 读 data[i] 使用 Shared Memory: Global Memory → Thread 0 读 data[i] → 存入 Shared Memory Shared Memory → Thread 1 读 data[i] ← 高速读取 Shared Memory → Thread 2 读 data[i]典型用法Tiled Matrix Multiplication详细解释可看【infra之路】Tiled Matrix Multiplication详解__global__ void matMulShared(float *A, float *B, float *C, int N) { // 每个 Block 负责 C 的一个 TILE_SIZE × TILE_SIZE 子矩阵 __shared__ float sA[TILE_SIZE][TILE_SIZE]; // 声明共享内存 __shared__ float sB[TILE_SIZE][TILE_SIZE]; int row threadIdx.y blockIdx.y * TILE_SIZE; int col threadIdx.x blockIdx.x * TILE_SIZE; float sum 0.0f; // 沿 K 维度分块每次加载一个 tile 到 shared memory for (int t 0; t N / TILE_SIZE; t) { // 协作加载每个线程负责加载 A 和 B 的一个元素 sA[threadIdx.y][threadIdx.x] A[row * N (t * TILE_SIZE threadIdx.x)]; sB[threadIdx.y][threadIdx.x] B[(t * TILE_SIZE threadIdx.y) * N col]; __syncthreads(); // 必须同步确保所有线程都加载完了 // 在 tile 内做计算 for (int k 0; k TILE_SIZE; k) { sum sA[threadIdx.y][k] * sB[k][threadIdx.x]; } __syncthreads(); // 必须同步确保所有线程都用完了再加载下一个 tile } C[row * N col] sum; }Bank Conflicts存储体冲突Shared Memory 内部被分成32 个 Bank与 warp 的 32 个线程对应。理想情况下每个线程访问不同 bank没有冲突。Bank 分配规则地址 addr 落在 bank (addr / 4) % 32 假设 4 字节元素 Bank 0: addr[0], addr[32], addr[64], ... Bank 1: addr[1], addr[33], addr[65], ... ... Bank 31: addr[31], addr[63], addr[95], ...冲突场景同一 warp 中多个线程访问同一 bank 的不同地址 → 串行化N 个线程冲突 → N-way bank conflict。// ✗ 所有线程访问同一列 → 全部命中 bank 0 → 32-way conflict sA[threadIdx.x][0] // 所有线程读 sA 的第 0 列 // ✓ 每个线程访问不同行不同列 → 无 conflict sA[threadIdx.x][threadIdx.x] // 对角线访问经典解决方案给二维 Shared Memory 数组加 padding// 原来32×32列访问会导致 bank conflict __shared__ float sA[32][32]; // 解决32×33错开 bank 映射 __shared__ float sA[32][33]; // 多一列 padding // 现在 sA[0][0]→bank 0, sA[1][0]→bank 1, sA[2][0]→bank 2...特殊情况所有线程访问同一地址broadcast不会产生 conflict硬件会广播该值给所有线程。Register寄存器寄存器是最快的存储每个线程私有。编译器会自动把局部变量分配到寄存器。关键权衡寄存器使用 vs Occupancy每个 SM 有固定数量的寄存器A100: 65536 个 32-bit每个线程使用的寄存器越多SM 能容纳的活跃线程数越少活跃线程数少 → 可用 warp 少 → 内存延迟隐藏不住 → 性能下降例A100每 SM 65536 个寄存器最大 2048 个线程 如果每线程用 32 个寄存器65536 / 32 2048 个线程 → 满 occupancy ✓ 如果每线程用 128 个寄存器65536 / 128 512 个线程 → 只有 25% occupancy ✗实践建议避免在 kernel 中使用大量局部变量或大数组会被溢出到 Local Memory即慢速 Global Memory用__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)提示编译器优化寄存器分配用nsight compute查看实际寄存器使用量综合示例Shared Memory 在 Attention 中的应用FlashAttention 的核心思想就是把本课讲的内存层次运用到极致标准 Attention 的问题 O(N²) 的 Attention 矩阵太大必须写到 Global MemoryHBM → 两次 HBM 读写写 S 矩阵 读 S 矩阵带宽瓶颈 FlashAttention 的解法 把 Q, K, V 切成小块tiles 每个 tile 加载到 SRAMShared Memory 在 SRAM 内完成局部 attention 计算 直接累加到输出不需要把整个 S 矩阵写回 HBM → HBM 访问量从 O(N²) 降到 O(N)这正是第 3-4 周要深入的 FlashAttention 原理但本质就是今天讲的把数据搬到更快的内存层级减少对慢速内存的访问。本课小结概念要点内存层次Register Shared L1 L2 Global速度差几个数量级Coalesced Access同一 warp 的线程访问连续地址 → 合并为 1 次内存事务矩阵访问threadIdx.x对应列方向内存连续方向才能合并Shared MemoryBlock 内共享程序员手动管理用__syncthreads()同步Bank Conflict多线程访问同 bank 不同地址会串行化用 padding 解决Broadcast所有线程读同一地址不冲突硬件自动广播Register 权衡每线程寄存器越多 → occupancy 越低 → 可能性能下降FlashAttention 本质把 attention 计算搬到 SRAM减少 HBM 访问自检问题在继续下一课之前确认你理解了这些一个 warp 访问 Global Memory 时32 个线程访问连续地址需要几次内存事务答1 次为什么操作矩阵时要让threadIdx.x对应列而不是行答列方向内存连续满足 coalesced access__syncthreads()在使用 Shared Memory 时为什么必须答确保所有线程都完成加载/消费避免数据竞争如果 kernel 中声明了一个float arr[100]的局部数组它会被放在哪答寄存器放不下时溢出到 Local Memory即 Global Memory很慢