你有没有遇到过这种情况写 CUDA kernel 时shared memory 用得很“标准”global memory 访问也尽量合并了但性能就是上不去。用 Nsight Compute 一分析发现 Shared Memory Bank Conflict 占了很大一块几十个线程都在抢同一个 bank访存效率被硬件强制串行化。这时候除了给 shared memory 加 padding还有一个更灵活的优化手段Swizzling。本文会从 bank conflict 的底层规则讲起拆解 XOR Swizzle 的原理并给出一套完整可运行的矩阵转置示例对比朴素实现、padding 实现和 swizzle 实现的性能差异。适合已经能写出简单 CUDA kernel、想进一步提升访存效率的开发者。1. 背景与核心概念1.1 Shared Memory 为什么快Shared memory 是 GPU 上一块可编程管理的片上存储器位于每个 SMStreaming Multiprocessor内部。它的访问延迟远低于 global memory通常只有几十个周期而 global memory 的延迟动辄几百个周期。因此很多 CUDA 优化策略都会把频繁访问的数据先搬进 shared memory让一个 block 内的线程反复复用这些数据。常见的应用场景包括矩阵乘法和矩阵转置。卷积神经网络的前向/反向传播。图像滤波、FFT、归约。shared memory 使用得当可以大幅减少对 global memory 的访问次数是 CUDA 性能优化里性价比很高的一步。1.2 什么是 Bank 和 Bank Conflictshared memory 在硬件上被划分为 32 个 Bank每个 Bank 的宽度通常是 4 字节。这意味着在一个 warp32 个线程同时访问 shared memory 时硬件会尽量让每个线程访问不同的 Bank从而实现一次指令周期内完成访问。如果 2 个或更多线程同时访问同一个 Bank 的不同地址就会发生 Bank Conflict。硬件只能把冲突的访问拆成多个周期依次执行性能随之下降。最坏情况下32 个线程都访问同一个 Bank效率只有理想情况的 1/32。需要注意多个线程访问同一个 Bank 的同一地址时硬件会通过广播机制处理不产生冲突。1.3 什么是 SwizzlingSwizzling 可以理解为一种“地址重排技术”。它不改变数据总量也不改变存储介质而是通过修改数据在 shared memory 中的布局映射关系让原本会落在同一 Bank 上的访问分散到不同 Bank。比较常见的是 XOR Swizzle它的核心公式在大多数教程里看起来很短swizzled_index original_index ^ (row mask);这个公式的使用范围却很讲究。如果用错了、用反了不但性能没提升结果还可能直接算错。2. 环境准备与版本说明本文示例使用 CUDA C 编写涉及 shared memory 和 kernel 性能分析。建议准备以下环境一块支持 CUDA 的 NVIDIA 显卡建议 Compute Capability 6.0 及以上。CUDA Toolkit 11.x 或更高版本本文代码没有使用新版本专属 API。支持 nvcc 编译的 Linux 或 Windows 环境。可选NVIDIA Nsight Compute用于分析 bank conflict。先检查本机 CUDA 环境nvidia-smi nvcc --versionnvidia-smi显示的是当前驱动支持的 CUDA 版本nvcc --version显示的是已安装的 CUDA Toolkit 版本。两者不一定相等但驱动版本最好不低于 Toolkit 要求的版本。如果你还没有配置好 CUDA 环境建议先完成安装并确认nvcc能被正常调用再继续本文的实验。不同项目的实际版本可能不同本文重点在于思路和代码结构编译参数需要根据本机 GPU 架构适当调整。3. Shared Memory 访问与 Swizzling 原理拆解3.1 Bank 索引是怎么算出来的以最常见的 4 字节数据例如float为例shared memory 的 bank 索引可以近似理解为bank (byte_offset / 4) % 32也就是说地址按 4 字节为一个粒度均匀映射到 32 个 bank 上。对一个二维 tile__shared__ float tile[TILE][TILE];如果按行主序存储那么tile[row][col]的地址是地址 row * TILE * 4 col * 4于是bank (row * TILE col) % 32当TILE是 32 的倍数时row * TILE对 32 取余为 0所以 bank 只由col决定。也就是说如果很多线程同时访问同一列的不同行它们会全部落到同一个 bank 上形成严重冲突。3.2 Padding 策略与局限Padding 是最经典的解法。把 tile 声明成__shared__ float tile[TILE][TILE 1];多出来的 1 列不需要存数据只是用来“错开”行与行之间的地址间隔。这时tile[row][col]的地址变成地址 row * (TILE 1) * 4 col * 4bank 变成bank (row * (TILE 1) col) % 32因为TILE 1通常不是 32 的倍数所以不同行相同列的元素会被分散到不同 bank。Padding 的优点是非常简单代码改动极小。缺点是它占用了额外的 shared memory并且对“行主序固定尺寸 tile”场景更有效对更复杂的不规则访问模式效果可能不够理想。3.3 XOR Swizzle 核心公式XOR Swizzle 的思路是在访问 shared memory 前对列索引做一次异或变换。最简单的形式是int swizzled_col threadIdx.x ^ threadIdx.y;其中threadIdx.x可以理解为列号threadIdx.y可以理解为行号。为什么这个操作能消除 bank conflict关键点在于固定threadIdx.y时threadIdx.x从 0 变到 31。threadIdx.x ^ threadIdx.y的结果仍然是 0 到 31 的一个排列。因此原本“同一 bank、不同线程”的情况会被打散到 32 个不同 bank。当 tile 尺寸是 64 或更大时最好对行号取掩码int swizzled_col threadIdx.x ^ (threadIdx.y (TILE - 1));