CUDA编程：Shared Memory Bank Conflict 与 Padding 优化

Michael阿明

发布于 2026-05-29 12:06:55

110

文章目录

1. 核心原理：Shared Memory 为什么也会冲突？
2. 用数组理解 bank 映射
3. 为什么矩阵转置容易出现 bank conflict？
4. 为什么 `tile[32][33]` 能解决问题？
5. 实验：矩阵转置 no-padding vs padding
6. 答疑
- 问题 1：为什么两个 kernel 结果都正确，但速度不同？
- 问题 2：为什么 `tile[32][32]` 容易冲突？
- 问题 3：为什么 `tile[32][33]` 有效？
7. 总结

Global Memory 的 Memory Coalescing：让 warp 内线程尽量访问连续 Global 内存地址。

但是也需要注意：

❝Shared Memory 很快，但如果访问方式不合理，也会因为 bank conflict 变慢。

摘要：

1. Shared Memory bank 是什么
2. 什么是 bank conflict
3. 为什么矩阵转置容易产生 bank conflict
4. 为什么 tile[32][33] 可以减少 bank conflict
5. 通过实验比较 no-padding 和 padding 的性能差异

1. 核心原理：Shared Memory 为什么也会冲突？

Shared Memory 它内部被分成多个 bank。

在很多 NVIDIA GPU 上，可以简化理解为：

Shared Memory ≈ 32 个 bank
Warp = 32 个线程

理想情况下，一个 warp 的 32 个线程同时访问 shared memory 时：

thread 0  -> bank 0
thread 1  -> bank 1
thread 2  -> bank 2
...
thread 31 -> bank 31

这时访问可以并行完成，速度很快。

但是如果多个线程访问同一个 bank：

thread 0  -> bank 0
thread 1  -> bank 0
thread 2  -> bank 0
...

那么这些访问会被串行化。

这就是：

❝Bank Conflict，共享内存 bank 冲突。

2. 用数组理解 bank 映射

假设 shared memory 中存的是 float，一个 float = 4 bytes。

可以粗略理解为：

bank_id = index % 32

例如：

shared[0]  -> bank 0
shared[1]  -> bank 1
shared[2]  -> bank 2
...
shared[31] -> bank 31
shared[32] -> bank 0
shared[33] -> bank 1

所以如果 warp 内线程访问：

shared[threadIdx.x]

那么：

thread 0 -> shared[0]  -> bank 0
thread 1 -> shared[1]  -> bank 1
...
thread31 -> shared[31] -> bank31

没有冲突。

但如果访问：

shared[threadIdx.x * 32]

那么：

thread 0 -> shared[0]    -> bank 0
thread 1 -> shared[32]   -> bank 0
thread 2 -> shared[64]   -> bank 0
...
thread31 -> shared[992]  -> bank 0

全部撞到 bank 0，性能会明显下降。

在这里插入图片描述

3. 为什么矩阵转置容易出现 bank conflict？

矩阵转置优化通常会这样做：

Global Memory 连续读
↓
Shared Memory 暂存 tile
↓
Shared Memory 中完成局部转置
↓
Global Memory 连续写

经典 shared memory tile：

__shared__ float tile[32][32];

写入 shared memory 时：

tile[threadIdx.y][threadIdx.x]

读取 shared memory 时为了转置，会变成：

tile[threadIdx.x][threadIdx.y]

问题出在第二种访问。

对于 tile[32][32]，每一行长度是 32。按列访问时，很容易出现：

thread 0 -> tile[0][same_col]
thread 1 -> tile[1][same_col]
thread 2 -> tile[2][same_col]
...

由于每一行跨度刚好是 32 个 float，不同线程访问的地址间隔刚好映射到同一个 bank。

于是发生严重 bank conflict。

在这里插入图片描述

4. 为什么 `tile[32][33]` 能解决问题？

经典优化是把：

__shared__ float tile[32][32];

改成：

__shared__ float tile[32][33];

也就是每一行多加 1 个 float。

这样每一行的跨度从：

变成：

bank 映射就会错开。

原来：

第 0 行起点 -> bank 0
第 1 行起点 -> bank 0
第 2 行起点 -> bank 0

加 padding 后：

第 0 行起点 -> bank 0
第 1 行起点 -> bank 1
第 2 行起点 -> bank 2
...

这样按列访问时，不同线程更容易落到不同 bank，冲突显著减少。

这就是： Padding 优化。

在这里插入图片描述

5. 实验：矩阵转置 no-padding vs padding

用矩阵转置来观察 bank conflict。

我们实现两个优化版 transpose kernel：

1. transpose_shared_no_padding
   使用 tile[32][32]
   可能有 shared memory bank conflict

2. transpose_shared_padding
   使用 tile[32][33]
   通过 padding 减少 bank conflict

两个版本 Global Memory 读写模式基本一致，主要差别在 shared memory 的列访问是否冲突。

代码如下：

#include <cuda_runtime.h>

#include <cmath>
#include <cstdlib>
#include <iomanip>
#include <iostream>
#include <vector>

#define CUDA_CHECK(call)                                                        \
    do {                                                                        \
        cudaError_t err = call;                                                 \
        if (err != cudaSuccess) {                                               \
            std::cerr << "CUDA Error: " << cudaGetErrorString(err)              \
                      << " at " << __FILE__ << ":" << __LINE__ << std::endl;    \
            std::exit(EXIT_FAILURE);                                            \
        }                                                                       \
    } while (0)

constexprint TILE_DIM = 32;
constexprint BLOCK_ROWS = 8;

/*
 * Shared Memory transpose without padding.
 *
 * tile[32][32] 在转置读取时容易产生 bank conflict。
 */
__global__ void transpose_shared_no_padding(const float* in,
                                            float* out,
                                            int width,
                                            int height) {
    __shared__ float tile[TILE_DIM][TILE_DIM];

    int x = blockIdx.x * TILE_DIM + threadIdx.x;
    int y = blockIdx.y * TILE_DIM + threadIdx.y;

    /*
     * 读取 input。
     *
     * threadIdx.x 连续，因此 Global Memory 读取是 coalesced 的。
     */
    for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS) {
        int yy = y + j;
        if (x < width && yy < height) {
            tile[threadIdx.y + j][threadIdx.x] = in[yy * width + x];
        }
    }

    __syncthreads();

    /*
     * 交换 blockIdx.x 和 blockIdx.y，写出转置后的 tile。
     */
    int out_x = blockIdx.y * TILE_DIM + threadIdx.x;
    int out_y = blockIdx.x * TILE_DIM + threadIdx.y;

    /*
     * 这里读取 tile[threadIdx.x][threadIdx.y + j]。
     * 对 tile[32][32] 来说，按列读 shared memory 容易发生 bank conflict。
     */
    for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS) {
        int oy = out_y + j;
        if (out_x < height && oy < width) {
            out[oy * height + out_x] = tile[threadIdx.x][threadIdx.y + j];
        }
    }
}

/*
 * Shared Memory transpose with padding.
 *
 * tile[32][33] 通过多加 1 列打散 bank 映射，减少 bank conflict。
 */
__global__ void transpose_shared_padding(const float* in,
                                         float* out,
                                         int width,
                                         int height) {
    __shared__ float tile[TILE_DIM][TILE_DIM + 1];

    int x = blockIdx.x * TILE_DIM + threadIdx.x;
    int y = blockIdx.y * TILE_DIM + threadIdx.y;

    /*
     * Global Memory 连续读取。
     */
    for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS) {
        int yy = y + j;
        if (x < width && yy < height) {
            tile[threadIdx.y + j][threadIdx.x] = in[yy * width + x];
        }
    }

    __syncthreads();

    int out_x = blockIdx.y * TILE_DIM + threadIdx.x;
    int out_y = blockIdx.x * TILE_DIM + threadIdx.y;

    /*
     * 仍然是 tile[threadIdx.x][threadIdx.y + j]，
     * 但由于每行长度是 33，bank 映射错开，冲突减少。
     */
    for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS) {
        int oy = out_y + j;
        if (out_x < height && oy < width) {
            out[oy * height + out_x] = tile[threadIdx.x][threadIdx.y + j];
        }
    }
}

/*
 * CPU reference transpose，用于结果校验。
 */
void transpose_cpu(const std::vector<float>& in,
                   std::vector<float>& out,
                   int width,
                   int height) {
    for (int y = 0; y < height; ++y) {
        for (int x = 0; x < width; ++x) {
            out[x * height + y] = in[y * width + x];
        }
    }
}

/*
 * 计时函数。
 *
 * 注意：
 * 这里测的是 kernel time，不包含 H2D/D2H。
 */
template <typename Launch>
float time_kernel(Launch launch, int repeat) {
    /*
     * warmup，避免第一次 kernel 启动影响正式计时。
     */
    launch();
    CUDA_CHECK(cudaDeviceSynchronize());

    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));

    float total_ms = 0.0f;

    for (int i = 0; i < repeat; ++i) {
        CUDA_CHECK(cudaEventRecord(start));

        launch();

        CUDA_CHECK(cudaEventRecord(stop));
        CUDA_CHECK(cudaEventSynchronize(stop));

        float ms = 0.0f;
        CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
        total_ms += ms;
    }

    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));

    return total_ms / repeat;
}

bool check_result(const std::vector<float>& ref,
                  const std::vector<float>& out,
                  float eps = 1e-5f) {
    if (ref.size() != out.size()) {
        returnfalse;
    }

    for (size_t i = 0; i < ref.size(); ++i) {
        float diff = std::fabs(ref[i] - out[i]);
        if (diff > eps) {
            std::cerr << "Mismatch at " << i
                      << ", ref=" << ref[i]
                      << ", out=" << out[i]
                      << ", diff=" << diff << std::endl;
            returnfalse;
        }
    }

    returntrue;
}

int main(int argc, char** argv) {
    int width = 4096;
    int height = 4096;
    int repeat = 10;

    if (argc >= 2) {
        width = std::atoi(argv[1]);
    }
    if (argc >= 3) {
        height = std::atoi(argv[2]);
    }
    if (argc >= 4) {
        repeat = std::atoi(argv[3]);
    }

    size_t count = static_cast<size_t>(width) * static_cast<size_t>(height);
    size_t bytes = count * sizeof(float);

    std::cout << "Shared Memory Bank Conflict\n";
    std::cout << "Matrix size : " << height << " x " << width << "\n";
    std::cout << "Data size   : " << bytes / 1024.0 / 1024.0 << " MB\n";
    std::cout << "Repeat      : " << repeat << "\n";

    std::vector<float> h_in(count);
    std::vector<float> h_ref(count);
    std::vector<float> h_no_padding(count);
    std::vector<float> h_padding(count);

    for (size_t i = 0; i < count; ++i) {
        h_in[i] = static_cast<float>((i * 17 + 13) % 1000) * 0.001f;
    }

    /*
     * CPU reference。
     * 大矩阵时会慢一些，但 4096x4096 可以接受。
     */
    transpose_cpu(h_in, h_ref, width, height);

    float* d_in = nullptr;
    float* d_out_no_padding = nullptr;
    float* d_out_padding = nullptr;

    CUDA_CHECK(cudaMalloc(&d_in, bytes));
    CUDA_CHECK(cudaMalloc(&d_out_no_padding, bytes));
    CUDA_CHECK(cudaMalloc(&d_out_padding, bytes));

    CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));

    dim3 block(TILE_DIM, BLOCK_ROWS);
    dim3 grid((width + TILE_DIM - 1) / TILE_DIM,
              (height + TILE_DIM - 1) / TILE_DIM);

    std::cout << "Block       : (" << block.x << ", " << block.y << ")\n";
    std::cout << "Grid        : (" << grid.x << ", " << grid.y << ")\n";

    auto launch_no_padding = [&]() {
        transpose_shared_no_padding<<<grid, block>>>(d_in, d_out_no_padding, width, height);
        CUDA_CHECK(cudaGetLastError());
    };

    auto launch_padding = [&]() {
        transpose_shared_padding<<<grid, block>>>(d_in, d_out_padding, width, height);
        CUDA_CHECK(cudaGetLastError());
    };

    float no_padding_ms = time_kernel(launch_no_padding, repeat);
    float padding_ms = time_kernel(launch_padding, repeat);

    CUDA_CHECK(cudaMemcpy(h_no_padding.data(),
                          d_out_no_padding,
                          bytes,
                          cudaMemcpyDeviceToHost));

    CUDA_CHECK(cudaMemcpy(h_padding.data(),
                          d_out_padding,
                          bytes,
                          cudaMemcpyDeviceToHost));

    bool ok_no_padding = check_result(h_ref, h_no_padding);
    bool ok_padding = check_result(h_ref, h_padding);

    /*
     * 转置大致是读一次 input + 写一次 output。
     */
    double moved_bytes = 2.0 * static_cast<double>(bytes);

    double no_padding_bw = moved_bytes / (no_padding_ms / 1000.0) / 1e9;
    double padding_bw = moved_bytes / (padding_ms / 1000.0) / 1e9;

    std::cout << std::fixed << std::setprecision(4);

    std::cout << "\n[Timing]\n";
    std::cout << "No-padding time : " << no_padding_ms << " ms\n";
    std::cout << "Padding time    : " << padding_ms << " ms\n";
    std::cout << "Speedup         : " << no_padding_ms / padding_ms << "x\n";

    std::cout << "\n[Effective Bandwidth]\n";
    std::cout << "No-padding BW   : " << no_padding_bw << " GB/s\n";
    std::cout << "Padding BW      : " << padding_bw << " GB/s\n";

    std::cout << "\n[Check]\n";
    std::cout << "No-padding check: " << (ok_no_padding ? "PASS" : "FAIL") << "\n";
    std::cout << "Padding check   : " << (ok_padding ? "PASS" : "FAIL") << "\n";

    CUDA_CHECK(cudaFree(d_in));
    CUDA_CHECK(cudaFree(d_out_no_padding));
    CUDA_CHECK(cudaFree(d_out_padding));

    return (ok_no_padding && ok_padding) ? 0 : 1;
}

看到输出结果：

CUDA Lesson 8: Shared Memory Bank Conflict
Matrix size : 4096 x 4096
Data size   : 64 MB
Repeat      : 10
Block       : (32, 8)
Grid        : (128, 128)

[Timing]
No-padding time : 1.2885 ms
Padding time    : 0.6721 ms
Speedup         : 1.9172x

[Effective Bandwidth]
No-padding BW   : 104.1683 GB/s
Padding BW      : 199.7107 GB/s

[Check]
No-padding check: PASS
Padding check   : PASS

6. 答疑

问题 1：为什么两个 kernel 结果都正确，但速度不同？

因为 no-padding 和 padding 只是 shared memory 布局不同，数学计算没有变。

所以结果应该一致。

但 shared memory 的 bank 映射不同，导致访问效率不同。

问题 2：为什么 `tile[32][32]` 容易冲突？

因为 32 刚好和 bank 数量相同。

按列访问时，相邻线程访问地址间隔是 32 个 float，容易映射到同一个 bank。

问题 3：为什么 `tile[32][33]` 有效？

因为每行多 1 个元素，让下一行起始位置错开一个 bank。

这样按列读取时，线程不再集中撞同一个 bank。

7. 总结

在这里插入图片描述

核心结论是：

Shared Memory 很快，但也有内部结构
Shared Memory 被划分成多个 bank
一个 warp 内多个线程访问同一个 bank，会发生 bank conflict
Bank conflict 会导致 shared memory 访问串行化
矩阵转置中的列访问非常容易触发 bank conflict
tile[32][33] 通过 padding 打散 bank 映射
Padding 代价很小，但可能明显提升性能

一句话总结：

❝Shared Memory 优化不仅是“用不用 shared memory”，还要看“怎么访问 shared memory”；padding 是解决矩阵转置 bank conflict 的经典技巧。

本文参与腾讯云自媒体同步曝光计划，分享自微信公众号。

原始发表：2026-05-26，如有侵权请联系 cloudcommunity@tencent.com 删除

优化

本文分享自 Michael阿明微信公众号，前往查看

如有侵权，请联系 cloudcommunity@tencent.com 删除。

本文参与腾讯云自媒体同步曝光计划，欢迎热爱写作的你一起参与！

登录后参与评论

0 条评论

热度