Skip to content

refactor: rename kTile→kBlock and simplify copy infra across 07-14 - #15

Merged
ArthurinRUC merged 1 commit into
mainfrom
refactor/kBlock-naming
May 13, 2026
Merged

ArthurinRUC merged 1 commit into
mainfrom
refactor/kBlock-naming

Conversation

@ArthurinRUC

Copy link
Copy Markdown
Owner

Unify naming and copy-tiling across the 07-14 demos so the spec layer reads consistently:

  • Rename template params and constants kTileM/N/K → kBlockM/N/K (07-14).
  • For 07-10 (non-TMA G2S/S2G copies), replace the kThreadsPerWarp / kTileM_Copy / kTileN_Copy / kAlignedCopyItems* machinery with a single kBlockK_Copy = cute::min(64, kBlockK)/8 (and kBlockN_Copy for files that copy C/O), and a fixed (1, 8) val layout. The thread layout is now uniformly K-major for A/B and N-major for C/O, matching the cp.async 128b atom.
  • 08 dynamic_mma: kShmSize switched to cute::max(A+B+C, O) so the no-cvt-precision path that only stages O has correct sizing.

Verified on remote H200 (SM90): all 12 drivers (07 swizzling, 08 dynamic_mma, 09 pipelining x2, 10 gemm_api x2, 11 tma_load_store, 12 tma_multicast_reduce, 13 warpgroup_mma x2, 14 warp_specialization x2) compile and produce numerically correct output.

Unify naming and copy-tiling across the 07-14 demos so the spec layer
reads consistently:

- Rename template params and constants kTileM/N/K → kBlockM/N/K (07-14).
- For 07-10 (non-TMA G2S/S2G copies), replace the kThreadsPerWarp /
  kTileM_Copy / kTileN_Copy / kAlignedCopyItems* machinery with a single
  kBlockK_Copy = cute::min(64, kBlockK)/8 (and kBlockN_Copy for files
  that copy C/O), and a fixed (1, 8) val layout. The thread layout is
  now uniformly K-major for A/B and N-major for C/O, matching the
  cp.async 128b atom.
- 08 dynamic_mma: kShmSize switched to cute::max(A+B+C, O) so the
  no-cvt-precision path that only stages O has correct sizing.

Verified on remote H200 (SM90): all 12 drivers (07 swizzling, 08
dynamic_mma, 09 pipelining x2, 10 gemm_api x2, 11 tma_load_store,
12 tma_multicast_reduce, 13 warpgroup_mma x2, 14 warp_specialization x2)
compile and produce numerically correct output.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
@ArthurinRUC
ArthurinRUC merged commit a8ee382 into main May 13, 2026
5 of 6 checks passed
@ArthurinRUC
ArthurinRUC deleted the refactor/kBlock-naming branch May 13, 2026 07:55
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant