3. 共享内存优化：Bank Conflict避免、Padding技术、动态共享内存分配

共享内存，说白了就是GPU上的一块“手边缓存”。

它比全局内存快得多，延迟能低一个数量级。但用好它，没那么简单。

我刚开始做量化回测引擎时，就栽在共享内存的bank conflict上。明明数据都在片上，性能就是上不去。后来一查，原来是访问模式不对。

3.1 什么是Bank Conflict？

共享内存被划分成32个bank（对应warp的32个线程）。每个bank的带宽是固定的。如果多个线程同时访问同一个bank的不同地址，就会发生冲突。

举个例子：

// 假设 shared float data[32][32];
// 线程 i 访问 data[i][j] —— 没问题，列访问
// 线程 i 访问 data[j][i] —— 出问题了，行访问

为什么会这样？因为共享内存的地址映射是交错的。相邻的32位地址会落在不同的bank上。但如果你按行访问，同一行的数据都在同一个bank里。32个线程同时访问同一行，那就是32路bank conflict。

3.2 避免Bank Conflict的实战技巧

我个人习惯，写共享内存代码前先画个访问模式图。你想想看，如果每个线程访问的地址偏移量是线程ID的线性函数，那大概率会有冲突。

技巧一：改变数据布局

把行访问改成列访问。或者用结构体数组（SoA）代替数组结构体（AoS）。

// 冲突版本：AoS
struct Particle { float x, y, z; };
shared Particle particles[256];
// 线程i访问 particles[i].x —— 32路冲突！

// 无冲突版本：SoA
shared float px[256], py[256], pz[256];
// 线程i访问 px[i] —— 无冲突

技巧二：使用Padding技术

这是我最常用的方法。在数组每行末尾加一个或多个空元素，让bank的映射错开。

// 原始：32x32矩阵，行访问有32路冲突
shared float data[32][32];

// Padding：每行加1个元素，变成33列
shared float data[32][33];
// 现在行访问只有2路冲突，甚至无冲突

为什么Padding有效？因为bank的映射是基于地址的。加了padding后，同一行的数据不再连续映射到同一个bank，而是分散开了。

3.3 动态共享内存分配

静态共享内存大小在编译时就固定了。但很多时候，我们需要根据运行时数据量来调整。

动态共享内存的用法很简单：

// 内核声明时用 extern
__global__ void myKernel() {
    extern __shared__ float sharedData[];
    // 使用 sharedData
}

// 启动内核时指定大小
myKernel<<<grid, block, 256 * sizeof(float)>>>();

注意几个坑：

• 动态共享内存只能声明为一维数组
• 多个动态数组需要手动计算偏移量
• 大小必须在启动时确定，不能在设备端修改

我曾经在写期权定价引擎时，需要同时存储多个维度的数据。动态共享内存让我能根据合约数量灵活分配，避免了内存浪费。

// 多数组动态共享内存示例
__global__ void optionPricing(int n) {
    extern __shared__ char sharedMem[];
    float *prices = (float*)sharedMem;
    float *strikes = (float*)&prices[n];
    float *vols = (float*)&strikes[n];
    // 现在有3个动态数组，每个n个元素
}

3.4 知识体系图

下面这张图总结了共享内存优化的核心逻辑：

3.5 避坑指南

我曾经在写一个矩阵转置内核时，用了动态共享内存，但忘记计算padding。结果转置后的数据全是错的。排查了半天才发现是bank conflict导致的数据覆盖。

从那以后，我养成了几个习惯：

• 写共享内存代码前，先在纸上画访问模式
• 用cuda-memcheck工具检查bank conflict
• 动态共享内存的偏移量计算，单独写个宏或函数

嗯，共享内存优化这块，说白了就是“让每个线程走自己的路，别挤在一起”。掌握了bank conflict的规律、Padding的用法、动态分配的技巧，你的GPU代码性能就能上一个台阶。

---

公众号:蓝海资料掘金营，微信deep3321