全局内存访问优化（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） 时。

共享内存的结构

Shared Memory 被划分为 32 个等宽的、可以并行访问的内存模块，称为 Banks（存储体）。

• 无冲突： 当一个 Warp 中的 32 个线程分别访问 32 个不同的 Banks 时，访问可以同时完成。

• 有冲突： 当多个线程试图访问同一个 Bank 但不同地址的数据时，访问就会发生序列化（排队），这就叫 Bank Conflict。

核心规则

• Bank 数量： 始终为 32 个（正好对应一个 Warp 的大小）。

• 映射方式： 地址 n 映射到 Bank = (n mod 32)。

• 特例（Broadcast）： 如果一个 Warp 内的所有线程访问同一个 Bank 里的同一个地址，会触发广播机制，这不属于冲突。

具体例子

[32][32] 的同一列是同一个 Bank

在内存中，二维数组 data[32][32] 其实是线性排列的（行优先存储）。

• 第一行：data[0][0] 到 data[0][31]（占据地址 0 到 31）

• 第二行：data[1][0] 到 data[1][31]（紧跟其后，占据地址 32 到 63）

现在看 Bank 0 里住着谁：

• 第 0 个元素：data[0][0] (Index 0)

• 第 32 个元素：data[1][0] (Index 32)

• 第 64 个元素：data[2][0] (Index 64)

• ...

• 第 32* n 个元素：data[n][0]

[16][16]呢？

元素索引	对应物理 Bank	备注
data[0][0] (Index 0)	Bank 0	第一行开始
data[0][15] (Index 15)	Bank 15	第一行结束
data[1][0] (Index 16)	Bank 16	第二行开始
data[1][15] (Index 31)	Bank 31	第二行结束
data[2][0] (Index 32)	Bank 0	第三行开始，回到了 Bank 0

避免 Bank Conflict 的方法

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

• 按行访问而非按列：

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

• 增加 padding（填充列）：

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

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

为什么多加一列就能解决？

• 原本 data[0][31] 在 Bank 31，data[1][0] 就在 Bank 0。

• 现在多了一个空位，data[0][32] 占了 Bank 0，那么 data[1][0] 就被挤到了 Bank 1。

这样，当你纵向访问 data[0][0], data[1][0], data[2][0]... 时：

• data[0][0] -> Bank 0

• data[1][0] -> Bank 1

• data[2][0] -> Bank 2

• 本来在同一列（同 Bank）的元素，被“错位”到了不同的 Bank 中，冲突瞬间消失！

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

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