性能首先是数据移动问题

GPU 的峰值算力增长通常快于片外内存带宽增长。许多 kernel 并非算不动，而是算术单元在等数据。优化的核心不是“shared memory 一定快”，而是：

减少需要跨越慢层级的数据量，让每次搬来的数据产生更多有效计算，并让搬运与计算尽可能重叠。

速度快、容量小、作用域窄
┌─────────────────────────────────────┐
│ 每线程 registers                    │
├─────────────────────────────────────┤
│ 每 block shared memory / 每 SM L1   │  片上；具体组合随架构变化
├─────────────────────────────────────┤
│ 全 device 共享 L2 cache             │
├─────────────────────────────────────┤
│ device memory：GDDR / HBM           │  片外显存
├─────────────────────────────────────┤
│ PCIe / NVLink / C2C                 │
├─────────────────────────────────────┤
│ host memory / 其他 GPU memory       │
└─────────────────────────────────────┘
速度慢、容量大、作用域宽

延迟没有一个跨代恒定的“周期数”，缓存命中、访问模式、并发请求和架构都会改变结果。比背数字更有用的是理解作用域、带宽、复用和事务数量。

各存储空间

合并访问：以 warp 的事务看问题

对现代 CUDA 架构，理解全局内存访问的实用方法是：一个 warp 的地址覆盖了多少个对齐的 32-byte segment。事务越少，搬来的无用字节越少。

float（4 B），一个 warp 读 32 个连续元素：

lane       0  1  2  ...  7 | 8 ... 15 | 16 ... 23 | 24 ... 31
address    连续 128 B
segment       32 B         |   32 B   |    32 B    |    32 B
结果：理想情况下由 4 个 32 B 事务覆盖

若每 lane 读取 a[lane * 32]：
地址跨越很大范围，可能需要接近每 lane 一个事务，大量带宽被浪费。

因此行主序二维数组通常让 threadIdx.x 映射到列：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
float v = a[y * pitch_elements + x];

要同时考虑：

• 连续性：warp 内地址尽量相邻；
• 对齐：起始地址和数据类型对齐；
• 利用率：请求的字节有多少真正参与计算；
• 次数：同一数据是否被多个 thread/block 重复从显存加载；
• 读写对称性：优化读取时不要让写回变成 stride access。

AoS、SoA 与向量化

struct Particle { float x, y, z, mass; };
Particle* particles;  // AoS

struct Particles { float *x, *y, *z, *mass; };
Particles particles;  // SoA

如果一个 kernel 只读取 x，SoA 让 warp 请求连续 float；AoS 会把不使用的 y/z/mass 一起拖进事务。若每次都消费完整且自然对齐的结构体，AoS 或 AoSoA 也可能合适。

向量化访问如 float4 可以减少动态指令数并表达 16-byte 访问，但前提是：

• 指针对 alignof(float4) 对齐；
• 长度尾部单独处理；
• 每个分量确实需要；
• 编译后的 load/store 和实际吞吐确有改善。

未经证明的 reinterpret_cast<float4*> 可能造成未对齐行为或错误，不是免费的加速开关。

Shared memory：显式管理的数据缓存

Shared memory 的价值主要有三种：

1. 复用：一个 global load 被 block 内多个线程使用；
2. 重排：global 端保持合并访问，在 shared 中变换行列或邻域；
3. 通信：归约、扫描、直方图等 block 内交换中间结果。

静态与动态分配：

__shared__ float tile[32][33];

extern __shared__ unsigned char storage[];
float* values = reinterpret_cast<float*>(storage);
int* flags = reinterpret_cast<int*>(values + blockDim.x);

动态 shared memory 的不同对象需要自己处理偏移与对齐。launch 的第三个参数给出总字节数。

Bank conflict

官方当前技术表中 shared memory 有 32 个 bank。warp 内多个线程在同一条指令中访问同一 bank 的不同地址，请求会被拆分；多个线程读同一地址则可广播，不属于普通冲突。

经典转置：

shared tile[32][32]

写入 tile[ty][tx]：相邻 tx → 相邻 bank，通常无冲突
读出 tile[tx][ty]：相邻 tx → 地址步长 32 words → 落到同一 bank

改为 tile[32][33]：每行多一列，下一行 bank 映射错开

template<int TILE = 32, int ROWS = 8>
__global__ void transpose(const float* in, float* out, int width, int height) {
  __shared__ float tile[TILE][TILE + 1];

  int x = blockIdx.x * TILE + threadIdx.x;
  int y = blockIdx.y * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += ROWS)
    if (x < width && y + j < height)
      tile[threadIdx.y + j][threadIdx.x] = in[(y + j) * width + x];

  __syncthreads();

  x = blockIdx.y * TILE + threadIdx.x;
  y = blockIdx.x * TILE + threadIdx.y;
  for (int j = 0; j < TILE; j += ROWS)
    if (x < height && y + j < width)
      out[(y + j) * height + x] = tile[threadIdx.x][threadIdx.y + j];
}

启动 dim3 block(32, 8)。每线程循环处理四个元素，保持读写合并；+1 padding 消除典型的转置 bank conflict。实际优化仍需用 bank conflict 指标验证。

Register、spill 与局部数组

编译器通常把自动标量变量放入寄存器，但以下情况容易进入 local memory：

• 寄存器不足导致 spill；
• 索引无法在编译期解析的局部数组；
• 大型局部对象；
• 调试或编译设置改变分配。

检查：

nvcc -O3 -lineinfo -Xptxas=-v kernel.cu

看到 lmem、spill load/store 后，不应机械地加 --maxrregcount。限制寄存器也许提高 occupancy，却可能制造更多片外流量。应比较：kernel 时间、active warps、register 数、local load/store 与 stall 原因。

Host 与 Device 之间的内存策略

Managed memory 的“一个指针”不代表“没有数据移动”。离散 GPU 上页面仍可能在 host/device 间迁移。常用引导：

cudaMemPrefetchAsync(ptr, bytes, device, stream);
cudaMemAdvise(ptr, bytes, cudaMemAdviseSetReadMostly, device);

Hints 是性能建议而非正确性同步。Full Unified Memory、HMM、ATS 与硬件一致性能力取决于操作系统、驱动和平台，必须按官方 Unified Memory 章节检查当前系统语义。

Pitched allocation 与二维数据

cudaMallocPitch 可以给每行加入满足设备访问要求的 padding：

float* d = nullptr;
size_t pitch_bytes = 0;
cudaMallocPitch(&d, &pitch_bytes, width * sizeof(float), height);

__global__ void k(float* base, size_t pitch_bytes, int w, int h) {
  int x = ...;
  int y = ...;
  if (x < w && y < h) {
    auto* row = reinterpret_cast<float*>(
        reinterpret_cast<char*>(base) + y * pitch_bytes);
    row[x] = ...;
  }
}

Pitch 是字节数，不一定等于 width * sizeof(T)，不要误当元素数。

异步拷贝：从“先搬完再算”到流水线

传统 tile 循环：

load tile 0 → wait → compute tile 0 → load tile 1 → wait → compute tile 1

多阶段流水线：

copy engine:  [load 0][load 1][load 2][load 3]
compute:              [calc 0][calc 1][calc 2]

Compute Capability 8.0 起存在硬件加速的 global→shared memcpy_async 路径；9.0 起 TMA 扩展到更大的多维批量异步搬运。可通过 Cooperative Groups、cuda::pipeline/cuda::barrier 或更底层 API 使用。正确性要求生产者等待空槽、消费者等待拷贝完成，并在复用缓冲区前确认旧计算已经结束。

带宽计算

理论带宽由显存时钟、总线宽度与传输方式决定，以产品官方规格为准。Kernel 的有效带宽：

$$BW_{effective}=\frac{Bytes_{read}+Bytes_{written}}{time}$$

Vector Add 对每个 float 理想读取 a,b 共 8 B，写 c 4 B，即约 12N 字节；若考虑 write allocate、缓存或额外流量，应结合 profiler 的 DRAM bytes 解释。

不要把 profiler 的峰值百分比、有效带宽和产品标称带宽混为同一个量。

内存优化决策树

DRAM 吞吐高、计算吞吐低？
├─ 是：memory-bound
│  ├─ global 事务利用率低 → 改布局/合并/对齐
│  ├─ 重复读取多         → shared/register tiling 或融合
│  ├─ host-device 拷贝多  → 批处理、pinned、异步流水线
│  └─ 必要字节已最少     → 压缩/量化/换算法
└─ 否：
   ├─ L1/L2 miss + long scoreboard → 提高局部性/预取/并发 warp
   ├─ shared bank conflict         → 改索引/padding
   ├─ local load/store             → 检查 spill/局部数组
   └─ 吞吐都低                     → 并行度、依赖、同步或 launch 问题

官方参考

• CUDA Programming Guide: Unified and System Memory
• CUDA C++ Best Practices: Memory Optimizations
• CUDA Programming Guide: Asynchronous Data Copies
• CUDA Programming Guide: Compute Capabilities