企业官网建设流程全解析

CUDA 开发中，共享内存（Shared Memory）是优化访存的重要手段，但它并不是一块简单的快速缓存。很多 kernel 明明已经把数据放进了共享内存，性能却不升反降，原因往往就是 bank conflict。Shared Memory Swizzling（共享内存重排）是一种通过调整数据在共享内存中的布局来消除 bank conflict 的优化手段。这篇文章从共享内存的物理结构讲起，解释 bank conflict 的产生原理，然后用一个矩阵转置程序对比普通布局、padding 和 XOR swizzle 三种写法，最后介绍如何用 Nsight Compute 定位冲突，以及设计 swizzle 布局时容易踩的坑。

1. 为什么共享内存访问会变慢：从 Bank 冲突说起

1.1 共享内存的物理结构

共享内存在 GPU 的每个 SM（Streaming Multiprocessor）内部，容量通常只有几十 KB 到上百 KB，但带宽远高于全局内存。它之所以快，是因为硬件把共享内存划分成了 32 个 bank，每个 bank 是独立的存储单元，每个时钟周期可以服务一个 word 的访问请求。这个 word 的大小一般是 4 字节，对应一个float或int。

当 warp 里的 32 个线程同时访问共享内存时，硬件事务逻辑会检查这 32 个地址落在哪些 bank 上。如果每个地址落在不同的 bank，那么这一次访问可以在一个事务内完成。如果多个地址落在同一个 bank，硬件就必须把这个请求拆成多个事务，每个事务服务其中一个地址，然后再合成结果。这个拆分的次数，就是冲突路数。

例如，线程tid访问s[tid]，地址从 0 到 31，分别落在 bank 0 到 bank 31，这是最理想的情况：

bank: 0 1 2 ... 31 地址: 0 1 2 ... 31

但如果线程tid访问s[tid * 2]，地址是 0, 2, 4, ..., 62，取模 32 之后得到的是 0, 2, 4, ..., 30, 0, 2, 4, ..., 30。每个 bank 上有两个线程访问，这就是 2-way bank conflict，实际需要两个事务才能完成这次访问。

如果 32 个线程访问同一个地址s[0]，硬件会启动广播机制，所有线程拿到同一个值，这种情况不算冲突，代价仍然是一个事务。所以判断是否冲突，关键看是否有多于一个线程访问同一个 bank 的不同地址。

1.2 为什么二维 tile 也逃不过冲突

在实际 kernel 中，共享内存很少被当成一维数组直接使用，更多是二维 tile，比如__shared__ float tile[32][32]。二维数组在内存里是行优先连续存储的，地址计算方式是：

address = row * row_stride + col

bank 编号就是address % 32。当行宽row_stride是 32 时：

bank = (row * 32 + col) % 32 = col

这意味着 bank 只由列索引决定，与行索引无关。如果 warp 内线程访问的列索引相同，那么无论行索引怎么变化，所有线程都会落在同一个 bank 上。

这正是很多二维 tile 程序性能低下的根源。看起来tile[row][col]很自然，但对硬件的 bank 分布来说，可能存在严重的访问集中。矩阵转置就是一个典型例子，后面会用完整代码演示。

1.3 冲突路数的代价估算

共享内存访问总延迟大致等于无冲突延迟乘以冲突路数。一个 32-way conflict 会把一次访问变成 32 次事务，共享内存的吞吐优势基本被抹平。

实际项目里不能只凭感觉判断，需要用性能分析工具统计共享内存的冲突次数。优化前先确认冲突存在，优化后再确认冲突被消除，这样每一步都有依据。

2. 理解 Swizzling：用重排布局消除 Bank 冲突

2.1 Swizzling 的基本思想

Swizzling 这个词在图形学里通常指对通道重新排列，比如把 RGBA 变成 BGRA。在 CUDA shared memory 场景里，它指对数据在共享内存中的物理位置做一次重排，使得原本会落在同一 bank 的访问被分散到不同 bank。

关键不是改变数据内容，而是改变数据存放的“列位置”。比如原本应该存到tile[row][col]的数据，现在存到tile[row][swizzle(row, col)]。只要swizzle函数是一个双射，即每个有效的(row, col)都映射到唯一的列索引，并且不同的(row, col)不会映射到同一个位置，数据就不会丢失。

为什么重排能消除冲突？因为冲突的本质是“多个线程的地址取模后落在同一 bank”。原本 bank 只由 col 决定，加入 row 的参与后，bank 变成 col 和 row 的某种组合。warp 内线程如果 col 相同但 row 不同，bank 也会不同，冲突就被打散。

2.2 Padding 是最简单的重排变体

在讲 XOR swizzle 之前，先看更常见的 padding 方案。把共享内存声明从tile[32][32]改成tile[32][33]，行宽多一个元素：

__shared__ float tile[32][33];

此时地址计算变成：

address = row * 33 + col bank = address % 32 = (row + col) % 32

因为33 % 32 == 1，行号每增加 1，bank 编号也增加 1。这样即使 warp 内线程的 col 相同，只要 row 不同，bank 也会分散开。

Padding 的实现成本极低，只需要改一行声明。缺点是浪费了一个元素的存储空间，并且对于某些复杂的访问模式，简单的单列 padding 不够灵活。比如行宽已经是 33 的倍数时，冲突又会出现。

2.3 XOR Swizzle：用异或打散 bank

Padding 的本质是把行号的信息叠加到列偏移上。XOR swizzle 也做类似的事，但用的是按位异或。一个常见的最小公式是：

#define SWIZZLE(col, row) ((col) ^ (row))

存储时不再写tile[row][col]，而是写tile[row][SWIZZLE(col, row)]。

为什么 XOR 有效？因为异或有一个特性：对于一个固定的row值，col ^ row在col取遍 0 到 31 时，结果仍然是 0 到 31 的一个排列。也就是说，warp 内如果 col 连续变化、row 固定，那么实际访问的 bank 编号恰好是 0 到 31 的乱序排列，每个 bank 只被访问一次，没有冲突。

需要注意col ^ row的结果必须落在有效列索引范围内。如果 tile 宽度是 32，col是 0..31，row最好也只取低 5 位参与异或，否则结果可能超出列范围。在 32x32 的 tile 里，行号和列号都在 0..31，所以可以直接用完整坐标异或。更通用的写法是：

#define SWIZZLE(col, row) ((col) ^ ((row) & 31))

这样无论 row 多大，结果都保持在 0..31。

XOR swizzle 不增加额外存储，适合对共享内存容量敏感的场景。它和 padding 并不是互斥的，有些高性能库会同时使用 padding 和 XOR swizzle，进一步打散特殊访问模式。

3. 实战：矩阵转置的三种写法与性能对比

矩阵转置非常适合用来演示 bank conflict。它既要用到二维 tile，又要让线程从共享内存的列方向读取数据，访问模式会直接暴露冲突问题。下面三个 kernel 分别对应普通布局、padding 和 XOR swizzle。

3.1 环境准备和编译命令

本文代码基于 CUDA C++，假设已经安装好 CUDA Toolkit，并有一块支持共享内存的 NVIDIA GPU。代码里使用 32x32 的 thread block，N 取 1024，便于演示。

编译命令：

nvcc -arch=sm_80 -O2 -o transpose transpose.cu

-arch=sm_80对应 Ampere 架构。如果你的 GPU 是其他架构，换成对应的 compute capability，比如sm_75或sm_90。编译成功后，程序会依次运行三个 kernel 并输出耗时。

3.2 普通布局：会产生 32 路冲突的基线版本

#include <cuda_runtime.h> #include <cstdio> #include <cstdlib> const int TILE = 32; __global__ void transpose_baseline(const float* in, float* out, int n) { int tx = threadIdx.x; int ty = threadIdx.y; __shared__ float tile[TILE][TILE]; int row_in = blockIdx.y * TILE + ty; int col_in = blockIdx.x * TILE + tx; // 写入阶段：tile[ty][tx]，warp 内 tx 连续变化，bank 不冲突 tile[ty][tx] = in[row_in * n + col_in]; __syncthreads(); // 读取阶段：tile[tx][ty]，warp 内 tx 连续变化，但列 ty 固定 // 地址 = tx * 32 + ty，bank = ty，所有线程访问同一个 bank int row_out = blockIdx.x * TILE + ty; int col_out = blockIdx.y * TILE + tx; out[row_out * n + col_out] = tile[tx][ty]; }

主函数启动方式如下，后面两个 kernel 使用相同的网格和线程配置：

int main() { const int N = 1024; size_t bytes = N * N * sizeof(float); float* h_in = (float*)malloc(bytes); float* h_out = (float*)malloc(bytes); for (int i = 0; i < N * N; i++) { h_in[i] = i % 7; } float* d_in; float* d_out; cudaMalloc(&d_in, bytes); cudaMalloc(&d_out, bytes); cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice); dim3 block(TILE, TILE); dim3 grid(N / TILE, N / TILE); cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start); transpose_baseline<<<grid, block>>>(d_in, d_out, N); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms = 0; cudaEventElapsedTime(&ms, start, stop); printf("baseline: %.3f ms\n", ms); cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost); cudaDeviceSynchronize(); cudaFree(d_in); cudaFree(d_out); free(h_in); free(h_out); return 0; }

这个版本在读取tile[tx][ty]时会产生 32-way bank conflict。因为 warp 内ty相同，tx从 0 变到 31，地址是tx * 32 + ty，取模 32 后全部等于ty。

3.3 Padding 版本：一行声明消除冲突

__global__ void transpose_padding(const float* in, float* out, int n) { int tx = threadIdx.x; int ty = threadIdx.y; // 行宽从 32 改成 33，额外的一列用于打破 bank 对齐 __shared__ float tile[TILE][TILE + 1]; int row_in = blockIdx.y * TILE + ty; int col_in = blockIdx.x * TILE + tx; tile[ty][tx] = in[row_in * n + col_in]; __syncthreads(); int row_out = blockIdx.x * TILE + ty; int col_out = blockIdx.y * TILE + tx; out[row_out * n + col_out] = tile[tx][ty]; }

这里只改了一个地方：__shared__ float tile[TILE][TILE + 1]。读取tile[tx][ty]时，地址变为tx * 33 + ty，bank 变为(tx + ty) % 32。warp 内tx从 0 变到 31，ty固定，所以 bank 依次变化，不再集中到同一个 bank。

3.4 XOR Swizzle 版本：不浪费额外存储

__global__ void transpose_swizzle(const float* in, float* out, int n) { int tx = threadIdx.x; int ty = threadIdx.y; __shared__ float tile[TILE][TILE]; int row_in = blockIdx.y * TILE + ty; int col_in = blockIdx.x * TILE + tx; // 存储时对列做重排 tile[ty][tx ^ ty] = in[row_in * n + col_in]; __syncthreads(); int row_out = blockIdx.x * TILE + ty; int col_out = blockIdx.y * TILE + tx; // 读取时使用对应的逆重排 out[row_out * n + col_out] = tile[tx][ty ^ tx]; }

这个版本的核心是读写两侧的列索引都加入异或。存储时，原位置(ty, tx)的数据放在列tx ^ ty；读取时，原位置(tx, ty)的数据要从列ty ^ tx取出。因为异或满足交换律，所以tx ^ ty和ty ^ tx是同一个值，写读一致。

从 bank 分布看，写入阶段地址是ty * 32 + (tx ^ ty)，bank 是tx ^ ty，warp 内tx连续变化时是排列，无冲突。读取阶段地址是tx * 32 + (ty ^ tx)，bank 是ty ^ tx，同样无冲突。

3.5 性能对比怎么看

三个 kernel 都完成了同样的矩阵转置，但共享内存访问行为不同。实际运行耗时不是固定的，跟 GPU 架构、编译优化、N 的大小都有关系。下面表格描述的是相对行为而不是精确数值：

在自己的机器上测试时，建议用cudaEvent记录耗时，并多次运行取稳定值。普通版本会明显慢于另外两个版本，padding 和 XOR 的差距通常很小。

注意：矩阵转置中，写共享内存阶段的冲突和读阶段不一定相同。这里普通版本冲突发生在读阶段，写阶段反而是无冲突的。分析时一定要把读写分开看。

4. 用 Nsight Compute 定位 Bank 冲突并验证优化效果

4.1 为什么需要性能分析工具

单纯比较耗时只能看到优化有没有效果，看不到冲突在哪。Nsight Compute 是 NVIDIA 提供的 kernel 性能分析工具，可以统计共享内存 bank conflict 的次数、访存吞吐、指令数等指标。优化 shared memory 访问之前，最好先用它确认冲突的来源和量级。

4.2 运行 Nsight Compute

编译可执行文件后，用ncu命令分析某个 kernel：

ncu --set full ./transpose

--set full会采集大量指标，输出耗时较长。如果只想看共享内存冲突，可以用--metrics指定指标：

ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum ./transpose

指标名称在不同 CUDA 版本下可能略有差异，如果命令报错，可以在 Nsight Compute 的图形界面里打开 kernel 的 Memory Workload Analysis 页面，查看 Shared Memory 部分的 Bank Conflicts 一行。

4.3 如何解读结果

对于普通版本，读共享内存的冲突计数会很高，理论上每个线程的读取都产生 32-way conflict，所以总次数约等于线程数乘以 32 的某种折算。padding 版本和 XOR swizzle 版本的冲突计数应该接近 0。

如果两个优化版本的冲突计数不为 0，优先检查以下几点：

这里要注意，ncu默认会重新运行 kernel，所以不需要在代码里做特殊处理。如果程序依赖命令行参数，可以放在ncu后边。

4.4 正确性验证

性能分析之前，先确保转置结果是正确的。可以用简单的方式验证：让每个元素等于row * N + col，转置后检查out[row * N + col] == col * N + row。也可以在 CPU 上写一个朴素转置对比。两个优化版本都没有改变数据的逻辑位置，只是在共享内存里的物理位置不同，所以输出应该和普通版本完全一致。

5. 设计 Swizzle 布局的通用方法和常见坑

5.1 先分析访问模式，再设计 Swizzle

不是所有 shared memory 程序都需要 swizzle。在动手之前，先把 warp 内 32 个线程访问共享内存的地址列出来，计算每个地址的 bank 编号。如果 bank 分布已经均匀，加 swizzle 只会增加额外计算。

设计 swizzle 的步骤可以归纳为：

1. 确定线程到坐标的映射方式，通常是(threadIdx.x, threadIdx.y)。
2. 确定每次共享内存访问的表达式，比如tile[A(tx,ty)][B(tx,ty)]。
3. 把地址表达式代入bank = address % 32，看冲突路数。
4. 如果某维固定导致 bank 集中，在该维上叠加另一维的信息，常用方法是 padding 或 XOR。
5. 用ncu验证优化前后冲突计数，并对比耗时。

对于 32x32 的 tile，col ^ row是一个安全有效的 swizzle。对于其他尺寸，要保证异或的结果仍然落在列范围内。一个通用做法是只取坐标的低 5 位参与异或，例如：

#define SWIZZLE_INDEX(col, row) ((col) ^ ((row) & 31))

这样结果始终在 0..31 内，不会越界。但要注意，如果 tile 列宽小于 32，需要进一步使用掩码，或者改用 padding。

5.2 不同访存粒度下的 bank 计算

前面例子访问的是float，一个线程一次访问 4 字节，正好对应一个 bank。如果访问double，每个数据占 8 字节，会跨越两个 bank。如果访问float4，占 16 字节，会跨越 4 个 bank。硬件会把这种访问拆成多个 phase，冲突分析要按 4 字节 word 重新计算。

一个常见的经验是：对于向量化访问，swizzle 公式需要作用在“向量索引”上，而不是字节地址上。例如存储float4 tile[32][8]，每行 8 个float4，每个线程访问一个float4，bank 计算要按float4的元素个数扩展。不能简单套用标量版本的col ^ row。

5.3 常见的坑和排查路径

坑 1：只改读取，不改写入

如果存储时用tile[ty][tx]，读取时却用tile[tx][ty ^ tx]，数据位置对不上，最终结果会错乱。swizzle 的核心是读写规则必须对称，要么都用同一套公式，要么保证逆映射正确。

坑 2：异或结果越界

col ^ row的结果可能大于 31。例如col = 0，row = 32，结果就是 32。这种情况会访问到下一行去，轻则数据错乱，重则触发非法访问。建议在 helper 函数里统一处理，并加注释说明 row 必须小于 32，或者使用row & 31。

坑 3：忽略了同步

共享内存是 block 内共享的，写入后必须__syncthreads()，否则其他线程读到的可能是旧数据。因为 swizzle 改变了物理位置，有些线程可能访问到其他线程刚写入的位置，数据依赖更强，漏同步的后果更隐蔽。

坑 4：对所有访问都套用同一个 swizzle

同一个 kernel 里可能有多种访问模式，比如一次读 tile 的一行，另一次读 tile 的一列。它们对 swizzle 的响应不同。不要默认一个公式解决所有问题，要分别分析。

坑 5：过早优化

如果 kernel 的瓶颈不在 shared memory，而在全局内存访问、计算量或同步开销，swizzle 并不能带来收益。先用 profiling 工具看整体瓶颈，再决定优化方向。

5.4 可以复用的检查清单

下面的清单适合在写或改 swizzle 逻辑时逐项确认：

1. 是否列出了一次 warp 访问的完整地址集合？
2. 是否计算了每个地址的 bank 编号？
3. swizzle 函数是否是双射？会不会出现两个坐标映射到同一个位置？
4. swizzle 后的列索引是否始终小于 tile 列宽？
5. 所有读取路径和写入路径是否使用一致的映射？
6. 写共享内存后是否调用了__syncthreads()？
7. 是否用ncu对比了优化前后的 bank conflict 计数？
8. 是否验证了 kernel 输出结果与 CPU 参考实现一致？
9. 是否在多个数据规模下测试过性能？
10. 是否确认 shared memory 不是当前瓶颈？

5.5 扩展方向

XOR swizzle 不止能用于矩阵转置。在卷积计算中，stencil 回读共享内存时经常出现跨行访问；在 FFT 中，radix 合并需要按不同步长读取；在reduce中，部分归约也会遇到 bank 冲突。这些场景都可以用同样的思路：把固定维度的 bank 集中问题，通过异或或 padding 打散。

更复杂的布局还有 8x8 或 16x16 的 tile swizzle，会同时考虑 L2 缓存 line 和 shared memory bank 的对齐。这些在 CUTLASS 等高性能计算库中很常见。实际工程里如果用了 CUTLASS，通常不需要手写 swizzle，因为它已经封装好了。

起步阶段，建议先把共享内存的 bank 机制吃透，再用 Nsight Compute 反复验证动手写的 swizzle。比记住一堆复杂公式更重要的，是建立一套“看访问 pattern -> 推导 bank 分布 -> 设计重排 -> 用工具确认”的排查方法。这也是 CUDA 优化里最值得沉淀的能力。