TMA：从 Global Memory 搬到 Shared Memory
一个 warp 中的一个 thread 提交 2D copy，TMA engine 搬运整个 tile，并可在写入时应用 swizzle
Swizzle
无
128B
Row offset
Col offset
Global Memory
16×128 fp16（每个 cell = 16B = 1×8 fp16）
TMA Engine
数据搬运引擎
→
cp.async.bulk.tensor.2d
+ SWIZZLE_128B
Shared Memory
8×8 sectors (swizzled)
▸
一个 thread 提交请求
，TMA engine 异步完成搬运
▸
写入时自动应用 swizzle
，减少目标访问的 bank conflict
▸
双向搬运：
load（GMEM→SMEM）和 store（SMEM→GMEM）