TMA: Global to Shared Memory
One thread issues a 2D copy; the TMA engine moves the full tile and can swizzle it as it writes
Swizzle
None
128B
Row offset
Col offset
Global Memory
16×128 fp16 (each cell = 16B = 1×8 fp16)
TMA Engine
Data Mover
→
cp.async.bulk.tensor.2d
+ SWIZZLE_128B
Shared Memory
8×8 sectors (swizzled)
▸
One thread submits the request;
the TMA engine performs the asynchronous copy
▸
Swizzle on arrival
reduces bank conflicts for the target access pattern
▸
Copies in both directions:
load (GMEM→SMEM) and store (SMEM→GMEM)