B-08 证明TMA 再强也只是「怎么搬」。若从第一天就把字段交错成 AoS、或把该合并的写拆成跨步引擎只能更高效地搬废字节。本章回到数据布局AoS↔SoA、以及 Transpose 这种「布局变换算子」——一次调整如何改变 sector 事务与有效带宽。TL;DR工程结论口径RTX 5090 /sm_120CUDA eventmedianuseful payload GB/s。完整表见docs/results/B-09_layout.md。布局决定事务形状同一逻辑读一个floatAoS 常让 warp 打到多 sectorSoA 让相邻线程读相邻地址。CC≥6.0 按32B sector计数Best Practices。该换 SoA少字段热读时红利最大——本机touch_fields1时 SoA 相对 AoS~13.6×190 vs 2576 useful GB/s。别盲目全量永久 SoA触达↑红利收窄——本机加速比13.6 → 10.7 → 6.2 → 1.80touch 1/2/4/8touch8仍 SoA 更快但已近 2× 量级热路径总整记录读写或宿主强绑 AoS 时先估转换成本。Transpose 用 SMEM tile本机 naive 仅 copy 的~39%tiled / pad 达~91% / ~92%pad 略优挂钩 B-02 bank。判停主看 useful GB/s 与加速比NCU 本机AoS32vs SoA4sectors/request理想 float≈4naive transpose 写侧32、tiled 读写皆4。1. 问题引擎之后还卡在「第一天怎么排」B-01 已讲清 sector / 合并物理B-07/B-08 讲清怎么异步搬。本章只回答两件事——都是「布局决定事务」只是时间点不同问题本章交付一句话字段级访问该 AoS 还是 SoAaos/soatouch_fields扫怎么存长期布局必须改行列布局时怎么付费copy/transpose_*对照必须改存时怎么付费算子内 / 离线变换边界避免重讲章节已覆盖本章不重复B-01sector、合并、向量化叙事不重写「什么是合并」长教程B-02SMEM bank / paddingtranspose pad 只挂钩一句B-08TMA / mbarrier不扫 AI、不写 descriptorB-10Checklist 汇总本章只产「布局症状→处方」原料左AoS相邻线程的同一字段相隔sizeof(Struct)→ 多 sector。右SoA字段数组连续 → 少事务、高 useful/transferred。2. 物理模型stride 就是事务形状记宽记录struct { float f[8]; }32B。线程i读第 0 个字段AoS: p[i].f[0] → 地址 stride 32B warp: 0,32,64,… → 往往 32×32B sector若每线程各落一扇区 SoA: x[i] → 地址 stride 4B warp: 连续 128B → 理想约 4×32B sector工程含义少数字段热读AoS 每线程常各落一 32B sector未用字段占满扇区 → useful/transferred 差。触达字段↑同一 particle 的多字段落在同一 32B sector里时后续字段多半吃 L2/扇区复用AoS 的 DRAM 惩罚被摊销 → 与 SoA 的差距通常收窄。这不是「AoS 变合并了」墙钟仍可能 SoA 更快。口径useful payload 触达字段的读写字节禁止用「整 struct 名义带宽」粉饰 AoS。解读时本机sweep的绝对 GB/s 受 L2 容量与热数据影响看加速比随touch_fields的形状勿和总线利用率混谈。3. 决策表何时换布局、何时做 Transpose信号建议热循环按字段扫、每次只用 1少数字段SoA或列式作设备侧主布局几乎总是整记录读写且宿主 API 绑 AoS可暂留 AoS或kernel 内局部转 SoA / views扩展阅读算法需要行↔列矩阵 / 卷积 im2col 类算子内 tiled transpose或一次性离线转置后常驻NCUsectors/request相对理想值明显偏高先修布局 / 索引映射再谈 TMA / pipeline已 compute-bound布局差被算力盖住仍建议修布局省功耗与给后续算子留余量但别指望巨大墙钟加速SoA ≠ 银弹全量、频繁的 AoS↔SoA 转换本身吃带宽。转换有成本时优先保证最热 kernel 的触达路径是合并友好的。4. Transpose 处方把跨步留在 SMEM上GMEM 读合、写跨步 → 带宽地板。下GMEM→SMEM合→ tile 内转置 → SMEM→GMEM合。配置行为期望copy同行读写上限参照transpose_naive读合、写跨步远低于 copytranspose_tiledSMEM 重排接近 copytranspose_padtile[32][33]消 bank细节见 B-02Harris 经典结论合并全局访问是大头pad 再补 bank。本章用同一叙事做可复现 binary不重讲 bank 全家桶。5. 实验怎么设计项路径代码examples/02_memory_optim/09_layout_transform.cu结果docs/results/B-09_sweep.csv/B-09_modes.csv/B-09_layout.md绘图python scripts/plot_b09_layout.py一条主命令布局主结论./bin/02_memory_optim_09_layout_transform--modesweep一条命令transpose 默认 touch1 的 layout 全表./bin/02_memory_optim_09_layout_transform--modemodesmode问题aos/soa固定touch_fields的 useful GB/ssweeptouch_fields∈{1,2,4,8}加速比是否随触达比例收窄copy/transpose_*naive 地板 vs tiled/pad vs copy 天花板证据优先级CUDA eventmedian→ useful GB/s / 加速比NCU sectors/request 可选旁证不要把 ncu 附着时自打印 ms 当结论。DO_NCU1bashexamples/02_memory_optim/09_profile_layout.sh ncu-only5.1 本机实测RTX 5090 / sm_120平台与完整表[docs/results/B-09_layout.md]。重画python scripts/plot_b09_layout.pyTouch-fields sweep主结论touch_fieldsaos_mssoa_msaos GB/ssoa GB/sSoA/AoS10.1770.013190257613.5920.1750.016384409610.6640.1700.02779049236.2380.3230.17983215011.80Transpose vs copydim4096verify OKmodemedian_msuseful GB/svs copycopy0.06919571.00×transpose_naive0.1757680.39×transpose_tiled0.07517790.91×transpose_pad0.07418030.92×怎么读少字段SoA~13.6×——布局红利最大。触达↑加速比收窄到~1.8×形状符合大纲假设未要求回到 1。tiled/pad 逼近 copynaive 是明确地板pad 相对 tiled 的增量远小于 naive→tiled。SoA 在 touch2/4 的极高 useful GB/s 含 L2 贡献——对比看加速比勿当 HBM 总线打满。5.2 NCU 旁证sectors/requestDO_NCU1bashexamples/02_memory_optim/09_profile_layout.sh ncu-only不要读 ncu 附着时程序自打印的 ms。明细docs/results/B-09_layout.md。配置ld sec/reqst sec/reqAoS touch1 / 83232SoA touch1 / 844transpose_naive432transpose_tiled44怎么读SoA 打到理想合并AoS 无论 touch 多少sec/req 都钉在 32加速比收窄 ≠ 合并变好。naive 是「读合写炸」tiled 把写拉回 4——与 event 带宽表互证。6. 工程边界硬件不限 sm_90合并是全架构问题。防 DCE字段更新必须写回 device 可见存储示例对触达字段1。AoS 实现陷阱禁止Particle local p[i]; …; p[i]local;这种整结构体赋值——编译器可能一次搬满 32B测不到字段跨步。矩阵边长示例要求--dim为 32 的倍数tile32。与 TMA描述符 swizzle / WGMMA 友好布局是 Module D 故事本章只保证 GMEM 事务形状正确。7. 扩展阅读不抢后续 ModuleColfax, Matrix Transpose in CUTLASSCuTe 抽象同一物理处方AoS→SoA data viewsarXiv:2405.12507、arXiv:2512.05516注解驱动布局变换CPE’25DOI:10.1002/cpe.70199DynaSOAr 机制段arXiv:1810.11765下一章B-10把 B-01B-09 收成「症状 → 证据 → 处方」一页 Checklist8. 工程 SOP 与常见误区建议流程用--mode sweep看 SoA/AoS 是否随touch_fields收窄低 touch 且加速比明显 1 → 设备侧主路径改 SoA或列式需要行列变换 →--mode modes看 tiled/pad 是否接近 copy有 NCU对比 aos vs soa 的 sectors/request修布局优先于上 TMA汇总进 B-10 checklist判停touch1仍 SoA≈AoS → 查是否整 struct 赋值、L2 是否把工作集盖住可加大--n、N 是否过小NCUAoS 仍 ~32 sec/req、SoA ~4 → 布局诊断成立本机已验证若两者都近 4 却墙钟差很大 → 查别的瓶颈tiled 远低于 copy → 查 dim 对齐、bank试transpose_pad、是否误用 naive 索引正确性检查失败 → 先修下标 / tile 映射再看带宽只有转换成本、没有热路径收益 → 别为 SoA 而 SoA高频误区把 Host AoS API 便利性直接当成 Device 最优布局用「名义 GB/s」含未触达字段为 AoS 辩护整 struct 读写却声称在测「单字段跨步」上了 TMA/pipeline 却不修跨步写B-08 误区⑦的落地版把 SMEM bank 教程复读一遍却不报 copy 对照9. 小结与下一章三句话少字段热读 → SoA用sweep看红利何时收窄。Transpose 用 SMEM tile目标是逼近 copy不是「换个下标」。布局先于引擎B-08 搬得再好也救不了第一天选错的 stride。下一章B-10做 Module B Checklist从症状到证据到处方的统一表。10. 参考文献官方文档CUDA C Best Practices Guide — Coalesced Access to Global MemoryCUDA Programming Guide — Writing CUDA SIMT KernelsCoalesced Global MemoryNsight Compute — Compute Triage Guidesectors/request工程教程Mark Harris, An Efficient Matrix Transpose in CUDA C/CColfax, Tutorial: Matrix Transpose in CUTLASSSemih Güreşçi, GPU Memory: AoS vs SoA文献形状非本机绝对数实证 / 前沿AoS→SoA viewsarXiv:2405.12507扩展 arXiv:2512.05516DynaSOArarXiv:1810.11765Eberhardt et al., GPUDranoCAV’17Annotation-guided AoS→SoACPE’25DOI:10.1002/cpe.70199
B-09. 数据布局（AoS/SoA/Transpose）：一次布局调整带来的事务变化
发布时间：2026/7/31 17:31:56
B-08 证明TMA 再强也只是「怎么搬」。若从第一天就把字段交错成 AoS、或把该合并的写拆成跨步引擎只能更高效地搬废字节。本章回到数据布局AoS↔SoA、以及 Transpose 这种「布局变换算子」——一次调整如何改变 sector 事务与有效带宽。TL;DR工程结论口径RTX 5090 /sm_120CUDA eventmedianuseful payload GB/s。完整表见docs/results/B-09_layout.md。布局决定事务形状同一逻辑读一个floatAoS 常让 warp 打到多 sectorSoA 让相邻线程读相邻地址。CC≥6.0 按32B sector计数Best Practices。该换 SoA少字段热读时红利最大——本机touch_fields1时 SoA 相对 AoS~13.6×190 vs 2576 useful GB/s。别盲目全量永久 SoA触达↑红利收窄——本机加速比13.6 → 10.7 → 6.2 → 1.80touch 1/2/4/8touch8仍 SoA 更快但已近 2× 量级热路径总整记录读写或宿主强绑 AoS 时先估转换成本。Transpose 用 SMEM tile本机 naive 仅 copy 的~39%tiled / pad 达~91% / ~92%pad 略优挂钩 B-02 bank。判停主看 useful GB/s 与加速比NCU 本机AoS32vs SoA4sectors/request理想 float≈4naive transpose 写侧32、tiled 读写皆4。1. 问题引擎之后还卡在「第一天怎么排」B-01 已讲清 sector / 合并物理B-07/B-08 讲清怎么异步搬。本章只回答两件事——都是「布局决定事务」只是时间点不同问题本章交付一句话字段级访问该 AoS 还是 SoAaos/soatouch_fields扫怎么存长期布局必须改行列布局时怎么付费copy/transpose_*对照必须改存时怎么付费算子内 / 离线变换边界避免重讲章节已覆盖本章不重复B-01sector、合并、向量化叙事不重写「什么是合并」长教程B-02SMEM bank / paddingtranspose pad 只挂钩一句B-08TMA / mbarrier不扫 AI、不写 descriptorB-10Checklist 汇总本章只产「布局症状→处方」原料左AoS相邻线程的同一字段相隔sizeof(Struct)→ 多 sector。右SoA字段数组连续 → 少事务、高 useful/transferred。2. 物理模型stride 就是事务形状记宽记录struct { float f[8]; }32B。线程i读第 0 个字段AoS: p[i].f[0] → 地址 stride 32B warp: 0,32,64,… → 往往 32×32B sector若每线程各落一扇区 SoA: x[i] → 地址 stride 4B warp: 连续 128B → 理想约 4×32B sector工程含义少数字段热读AoS 每线程常各落一 32B sector未用字段占满扇区 → useful/transferred 差。触达字段↑同一 particle 的多字段落在同一 32B sector里时后续字段多半吃 L2/扇区复用AoS 的 DRAM 惩罚被摊销 → 与 SoA 的差距通常收窄。这不是「AoS 变合并了」墙钟仍可能 SoA 更快。口径useful payload 触达字段的读写字节禁止用「整 struct 名义带宽」粉饰 AoS。解读时本机sweep的绝对 GB/s 受 L2 容量与热数据影响看加速比随touch_fields的形状勿和总线利用率混谈。3. 决策表何时换布局、何时做 Transpose信号建议热循环按字段扫、每次只用 1少数字段SoA或列式作设备侧主布局几乎总是整记录读写且宿主 API 绑 AoS可暂留 AoS或kernel 内局部转 SoA / views扩展阅读算法需要行↔列矩阵 / 卷积 im2col 类算子内 tiled transpose或一次性离线转置后常驻NCUsectors/request相对理想值明显偏高先修布局 / 索引映射再谈 TMA / pipeline已 compute-bound布局差被算力盖住仍建议修布局省功耗与给后续算子留余量但别指望巨大墙钟加速SoA ≠ 银弹全量、频繁的 AoS↔SoA 转换本身吃带宽。转换有成本时优先保证最热 kernel 的触达路径是合并友好的。4. Transpose 处方把跨步留在 SMEM上GMEM 读合、写跨步 → 带宽地板。下GMEM→SMEM合→ tile 内转置 → SMEM→GMEM合。配置行为期望copy同行读写上限参照transpose_naive读合、写跨步远低于 copytranspose_tiledSMEM 重排接近 copytranspose_padtile[32][33]消 bank细节见 B-02Harris 经典结论合并全局访问是大头pad 再补 bank。本章用同一叙事做可复现 binary不重讲 bank 全家桶。5. 实验怎么设计项路径代码examples/02_memory_optim/09_layout_transform.cu结果docs/results/B-09_sweep.csv/B-09_modes.csv/B-09_layout.md绘图python scripts/plot_b09_layout.py一条主命令布局主结论./bin/02_memory_optim_09_layout_transform--modesweep一条命令transpose 默认 touch1 的 layout 全表./bin/02_memory_optim_09_layout_transform--modemodesmode问题aos/soa固定touch_fields的 useful GB/ssweeptouch_fields∈{1,2,4,8}加速比是否随触达比例收窄copy/transpose_*naive 地板 vs tiled/pad vs copy 天花板证据优先级CUDA eventmedian→ useful GB/s / 加速比NCU sectors/request 可选旁证不要把 ncu 附着时自打印 ms 当结论。DO_NCU1bashexamples/02_memory_optim/09_profile_layout.sh ncu-only5.1 本机实测RTX 5090 / sm_120平台与完整表[docs/results/B-09_layout.md]。重画python scripts/plot_b09_layout.pyTouch-fields sweep主结论touch_fieldsaos_mssoa_msaos GB/ssoa GB/sSoA/AoS10.1770.013190257613.5920.1750.016384409610.6640.1700.02779049236.2380.3230.17983215011.80Transpose vs copydim4096verify OKmodemedian_msuseful GB/svs copycopy0.06919571.00×transpose_naive0.1757680.39×transpose_tiled0.07517790.91×transpose_pad0.07418030.92×怎么读少字段SoA~13.6×——布局红利最大。触达↑加速比收窄到~1.8×形状符合大纲假设未要求回到 1。tiled/pad 逼近 copynaive 是明确地板pad 相对 tiled 的增量远小于 naive→tiled。SoA 在 touch2/4 的极高 useful GB/s 含 L2 贡献——对比看加速比勿当 HBM 总线打满。5.2 NCU 旁证sectors/requestDO_NCU1bashexamples/02_memory_optim/09_profile_layout.sh ncu-only不要读 ncu 附着时程序自打印的 ms。明细docs/results/B-09_layout.md。配置ld sec/reqst sec/reqAoS touch1 / 83232SoA touch1 / 844transpose_naive432transpose_tiled44怎么读SoA 打到理想合并AoS 无论 touch 多少sec/req 都钉在 32加速比收窄 ≠ 合并变好。naive 是「读合写炸」tiled 把写拉回 4——与 event 带宽表互证。6. 工程边界硬件不限 sm_90合并是全架构问题。防 DCE字段更新必须写回 device 可见存储示例对触达字段1。AoS 实现陷阱禁止Particle local p[i]; …; p[i]local;这种整结构体赋值——编译器可能一次搬满 32B测不到字段跨步。矩阵边长示例要求--dim为 32 的倍数tile32。与 TMA描述符 swizzle / WGMMA 友好布局是 Module D 故事本章只保证 GMEM 事务形状正确。7. 扩展阅读不抢后续 ModuleColfax, Matrix Transpose in CUTLASSCuTe 抽象同一物理处方AoS→SoA data viewsarXiv:2405.12507、arXiv:2512.05516注解驱动布局变换CPE’25DOI:10.1002/cpe.70199DynaSOAr 机制段arXiv:1810.11765下一章B-10把 B-01B-09 收成「症状 → 证据 → 处方」一页 Checklist8. 工程 SOP 与常见误区建议流程用--mode sweep看 SoA/AoS 是否随touch_fields收窄低 touch 且加速比明显 1 → 设备侧主路径改 SoA或列式需要行列变换 →--mode modes看 tiled/pad 是否接近 copy有 NCU对比 aos vs soa 的 sectors/request修布局优先于上 TMA汇总进 B-10 checklist判停touch1仍 SoA≈AoS → 查是否整 struct 赋值、L2 是否把工作集盖住可加大--n、N 是否过小NCUAoS 仍 ~32 sec/req、SoA ~4 → 布局诊断成立本机已验证若两者都近 4 却墙钟差很大 → 查别的瓶颈tiled 远低于 copy → 查 dim 对齐、bank试transpose_pad、是否误用 naive 索引正确性检查失败 → 先修下标 / tile 映射再看带宽只有转换成本、没有热路径收益 → 别为 SoA 而 SoA高频误区把 Host AoS API 便利性直接当成 Device 最优布局用「名义 GB/s」含未触达字段为 AoS 辩护整 struct 读写却声称在测「单字段跨步」上了 TMA/pipeline 却不修跨步写B-08 误区⑦的落地版把 SMEM bank 教程复读一遍却不报 copy 对照9. 小结与下一章三句话少字段热读 → SoA用sweep看红利何时收窄。Transpose 用 SMEM tile目标是逼近 copy不是「换个下标」。布局先于引擎B-08 搬得再好也救不了第一天选错的 stride。下一章B-10做 Module B Checklist从症状到证据到处方的统一表。10. 参考文献官方文档CUDA C Best Practices Guide — Coalesced Access to Global MemoryCUDA Programming Guide — Writing CUDA SIMT KernelsCoalesced Global MemoryNsight Compute — Compute Triage Guidesectors/request工程教程Mark Harris, An Efficient Matrix Transpose in CUDA C/CColfax, Tutorial: Matrix Transpose in CUTLASSSemih Güreşçi, GPU Memory: AoS vs SoA文献形状非本机绝对数实证 / 前沿AoS→SoA viewsarXiv:2405.12507扩展 arXiv:2512.05516DynaSOArarXiv:1810.11765Eberhardt et al., GPUDranoCAV’17Annotation-guided AoS→SoACPE’25DOI:10.1002/cpe.70199
