<<<CUDA C++ 学习路线grid · block · warp · lane
EN中文
阶段 2 · 执行模型与内存06 / 15约 22 分钟cuda/07_transpose.cu

内存合并：同一段代码，5 倍性能差距

warp 如何发起访存事务、什么是 stride 访问、AoS 与 SoA 的抉择

学完这一课你会
  • 理解一个 warp 的访存是如何被硬件合并成 32B sector 的
  • 识别代码中破坏合并的常见写法
  • 掌握 AoS → SoA 这个最有效的数据布局改造
  • 会用 shared memory 修复无法避免的非合并访问

GPU 访问显存不是按字节，而是按 32 字节的 sector 为单位。当一个 warp 的 32 个线程同时发起访存，硬件会把它们的地址合并成尽可能少的 sector 请求。理想情况下，32 个线程读连续的 32 个 float（128 字节）= 4 个 sector，一次搞定，零浪费。最坏情况下，32 个线程各读一个相距很远的 float = 32 个 sector = 1024 字节被搬进来，而你只用了其中 128 字节，87.5% 的带宽被扔掉了。

合并访问   c[i] = a[i]              stride-8 访问   c[i] = a[i * 8]

地址: 0  4  8 ... 124                  地址: 0  32  64 ... 992
线程: t0 t1 t2 ... t31                 线程: t0 t1  t2 ... t31

 ┌────┬────┬────┬────┐                 ┌────┐    ┌────┐    ┌────┐
 │sec0│sec1│sec2│sec3│                 │sec0│... │sec8│... │sec31│
 └────┴────┴────┴────┘                 └────┘    └────┘    └────┘
   全部命中，无浪费                        32 个 sector，每个只用 4/32 字节

 搬运 128 B / 使用 128 B = 100%         搬运 1024 B / 使用 128 B = 12.5%
合并 vs 跨步：同样读 32 个 float
▸打开内存合并模拟器调节 stride、offset 和元素宽度，实时看到一个 warp 会产生多少个 sector 事务、带宽利用率是多少。把 stride 从 1 拖到 2 再拖到 32，效率的塌陷非常直观。

破坏合并的四种典型写法

1. 跨步访问
1// ❌ 每个线程跳着取，warp 覆盖 32 * 8 * 4 = 1024 字节2__global__ void strided(const float* in, float* out, int stride) {3    int i = blockIdx.x * blockDim.x + threadIdx.x;4    out[i] = in[i * stride];5}6 7// ✅ 若 stride 来自「多通道数据」，改成让通道走 blockIdx，元素走 threadIdx8__global__ void byChannel(const float* in, float* out, int n) {9    int i = blockIdx.x * blockDim.x + threadIdx.x;   // 元素维，连续10    int c = blockIdx.y;                              // 通道维，warp 内不变11    out[c * n + i] = in[c * n + i];12}
2. 行列写反（二维索引最常见的错误）
1// ❌ threadIdx.x 对应行：warp 内 32 个线程各跨一整行2int row = blockIdx.x * blockDim.x + threadIdx.x;3int col = blockIdx.y * blockDim.y + threadIdx.y;4out[row * cols + col] = in[row * cols + col];5 6// ✅ threadIdx.x 对应列：warp 覆盖连续的 128 字节7int col = blockIdx.x * blockDim.x + threadIdx.x;8int row = blockIdx.y * blockDim.y + threadIdx.y;9out[row * cols + col] = in[row * cols + col];

3. AoS vs SoA：粒子系统的经典案例

结构体数组（AoS）——直觉上自然，性能上灾难
1struct Particle { float x, y, z, vx, vy, vz; };   // 24 字节2 3__global__ void stepAoS(Particle* p, float dt, int n) {4    int i = blockIdx.x * blockDim.x + threadIdx.x;5    if (i >= n) return;6    p[i].x += p[i].vx * dt;      // 相邻线程访问的 x 相距 24 字节7    p[i].y += p[i].vy * dt;8    p[i].z += p[i].vz * dt;9}10// 一个 warp 读 x 时跨越 32 * 24 = 768 字节，11// 且 24 不是 32 的因数，sector 还会错位，效率进一步下降。
数组结构体（SoA）——完美合并
1struct Particles {          // 每个字段是一个独立的连续数组2    float *x, *y, *z;3    float *vx, *vy, *vz;4};5 6__global__ void stepSoA(Particles p, float dt, int n) {7    int i = blockIdx.x * blockDim.x + threadIdx.x;8    if (i >= n) return;9    p.x[i] += p.vx[i] * dt;      // warp 读 x[0..31] = 连续 128 字节10    p.y[i] += p.vy[i] * dt;11    p.z[i] += p.vz[i] * dt;12}

4. 无法避免的非合并：用共享内存中转

矩阵转置是个典型例子：读和写这两个方向，你至多只能保住一个是合并的。解法是让全局内存的读写都保持合并，把「转置」这个乱序动作放到共享内存里完成——共享内存的随机访问代价远低于显存。

07_transpose.cu — 共享内存版转置
1#define TILE 322__global__ void transposeShared(const float* in, float* out, int w, int h) {3    __shared__ float tile[TILE][TILE + 1];   // +1 消除 bank conflict，下一课细讲4 5    int x = blockIdx.x * TILE + threadIdx.x;6    int y = blockIdx.y * TILE + threadIdx.y;7 8    // 读：warp 沿 x 连续 → 合并9    if (x < w && y < h) {10        tile[threadIdx.y][threadIdx.x] = in[y * w + x];11    }12    __syncthreads();13 14    // 写：交换 block 索引，让输出侧的 warp 也沿连续方向走 → 同样合并15    x = blockIdx.y * TILE + threadIdx.x;16    y = blockIdx.x * TILE + threadIdx.y;17    if (x < h && y < w) {18        out[y * h + x] = tile[threadIdx.x][threadIdx.y];   // 转置发生在共享内存里19    }20}

自测

先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。