全局内存访问优化（Coalesced Access）

什么是 Coalesced Access？

定义：一个 warp（32 个线程）在同一指令中访问全局内存时，如果这些访问请求可以合并成尽可能少的内存事务（通常是 32、64 或 128 字节对齐的块），就叫 coalesced。

条件:一个 warp 的线程访问 连续且对齐 的地址。

int tid = threadIdx.x + blockIdx.x * blockDim.x;
float val = d_array[tid];  // ✅ 连续访问 → Coalesced

优化技巧

以结构体对齐内存：使用 __align__(16) 或 float4。

为什么需要对齐？

• GPU 内存总线要求访问按 32/64/128 字节对齐，这样才能合并成一次事务。

• 如果内存对齐不好，warp 访问会拆分成多个事务，带宽利用率降低。

技巧

• 使用 float4（4 个 float 一起）保证 16 字节对齐。

• 或者使用 CUDA 对齐修饰符：

float4 vs float

float

• 单个 4 字节（32 bit）的浮点数。

• 每个线程访问 1 个 float 时，如果 warp 中 32 个线程访问地址连续（0,1,2,3...），CUDA 会把它们合并成1~2 个内存事务，性能好。

float4

• CUDA 提供的 矢量类型，表示 4 个连续的 float（总共 16 字节）。

• 优点：

  • 天然 16 字节对齐（满足 GPU 内存事务对齐要求）。

  • 每个线程一次加载 4 个浮点数，减少指令数，提高带宽利用率。

//设置float4
float4 f;
f.x = 1.1, f.y = 2.2, f.z = 3.3, f.w = 4.4;
float4 v = data[idx];  // 读取 4 个 float
result = v.x + v.y + v.z + v.w;

如果要读取大数组，使用 float4 可以让 每个线程批量读取，提高 coalesced 访问效率。

对比：

• 32 个线程一次访问 float → 128 字节（32*4）

• 32 个线程一次访问 float4 → 512 字节（32*16），如果对齐良好，GPU 可以用更少的事务完成。

__align__(n) 关键字

作用

• 强制结构体或变量的起始地址对齐到 n 字节边界。

• 为什么？因为 GPU（和 CPU）要求数据按一定字节对齐访问，否则：

  • 拆分访问 → 多次内存事务 → 性能差

  • 未对齐访问 → 有的设备直接报错

struct __align__(16) MyStruct {float x, y, z, w;
}; // 占 16 字节，起始地址必须是 16 的倍数

• 如果不加 __align__(16)，可能被编译器按 4 字节对齐排布，不符合 GPU 要求。

为什么 CUDA 推荐使用 float4 + 对齐？

• 全局内存的访问规则：按 32/64/128 字节事务合并。

• 如果 warp 32 线程访问 float（每个 4 字节），正好 128 字节，可以合并。

• 如果 warp 32 线程访问 float4（每个 16 字节），正好 512 字节，GPU 需要 4 个事务，但每个事务更大，吞吐率更高。

• 重要：必须保证起始地址按 float4 对齐，否则性能下降。

行优先存储，避免跨行访问

• CUDA 全局内存是按一维线性存储的，如果访问跨行，会破坏 coalesced。

• 例如，二维矩阵 A[M][N]，默认按行优先（row-major）存储：

内存布局: A[0][0], A[0][1], ..., A[0][N-1], A[1][0], ...

错误访问模式（列遍历）：

val = A[col][row]; // 每个线程跨 stride 访问

每个线程 stride 大，warp 访问不连续，性能差。

优化

• 保证 threadIdx.x 对应 最快变化维度（行访问），这样 warp 线程连续访问。

调换索引顺序，确保 threadIdx.x 是最快变化维度

• 原则：warp 线程访问地址必须连续。

• 如果你的算法天然是列操作，可以调整线程分布：

如果是行优先，那么变化最快的其实是列下标，threadIdx.x对应的也应该是列。如果算法要求列优先，可以对row和col进行调换

以下是行优先情况下col和row的写法

int col = blockIdx.x * blockDim.x + threadIdx.x;  // x 对应列
int row = blockIdx.y * blockDim.y + threadIdx.y;  // y 对应行

使用 Shared Memory 缓存 tile

为什么？

• 全局内存访问延迟大（400~600 cycles），共享内存延迟低（≈100x 更快）。

• 如果每个线程直接从全局内存多次访问，会拖慢性能。

• 解决：把要用的数据块（tile）加载到共享内存，所有线程复用，减少全局访问。

例子见上一篇文章：cuda编程笔记（9）--使用 Shared Memory 实现 tiled GEMM -CSDN博客

Bank Conflict

Bank Conflict（共享内存银行冲突） 是 CUDA 编程中的一个性能问题，发生在多个线程同时访问 共享内存（Shared Memory） 时。

共享内存的结构

• CUDA 的 共享内存被划分成多个 Bank，类似一个并行访问的“多路存储器”。

• 每个 Bank 可以在一个时钟周期内处理 1 个 32-bit 访问请求。

• Warp（32 个线程）同时访问共享内存时：

  • 如果 32 个线程访问 32 个不同的 Bank → 无冲突（完美并行）。

  • 如果 多个线程访问同一个 Bank 的不同地址 → 发生 Bank Conflict，访问会被 串行化，性能大幅下降。

具体原理

假设：

• 共享内存被分为 32 个 Bank

• 每个 Bank 宽度 = 4 字节（一个 float）

• 地址映射公式：

bank_id = (address_in_bytes / 4) % 32

例子

__shared__ float s[32][32];
按行优先存储（Row-major）：

• s[i][j] 的地址 = base + (i * 32 + j) * 4

情况 1：访问同一列

如果每个线程访问 s[threadIdx.x][k]（同一列 k），

• 地址 = base + (threadIdx.x * 32 + k) * 4

• bank_id = (threadIdx.x * 32 + k) % 32 = k（因为 threadIdx.x * 32 是 32 的倍数）

• 所有线程访问同一 Bank（k） → 严重冲突

情况 2：访问同一行

如果每个线程访问 s[k][threadIdx.x]（同一行 k），

• 地址 = base + (k * 32 + threadIdx.x) * 4

• bank_id = (k * 32 + threadIdx.x) % 32 = threadIdx.x

• 每个线程访问不同 Bank → 无冲突

避免 Bank Conflict 的方法

核心原则：让 warp 内的 32 个线程访问的地址尽量分布到不同的 bank。

• 按行访问而非按列：

  • 推荐：s[threadIdx.y][threadIdx.x]（X 对应列，变化最快）

• 增加 padding（填充列）：

  • 如果二维数组导致 bank 冲突，可以在第二维加一个“dummy 列”，让 stride ≠ 32：

__shared__ float s[TILE_SIZE][TILE_SIZE + 1];

使用结构化数据（float4）或 align：

• 一次加载多个元素，减少 warp 的 bank 竞争。