Skip to content

# docs: add layout algebra comment to make_tiled_copy_V_interleave_trans (#44) - #86

Open
XFDG wants to merge 1 commit into
Tencent:mainfrom
XFDG:doc-issue-44
Open

# docs: add layout algebra comment to make_tiled_copy_V_interleave_trans (#44)#86
XFDG wants to merge 1 commit into
Tencent:mainfrom
XFDG:doc-issue-44

Conversation

@XFDG

@XFDG XFDG commented Aug 24, 2026

Copy link
Copy Markdown

Summary

  • Add a detailed documentation comment to make_tiled_copy_V_interleave_trans() in src/attention/decode/sm90/util_kernels.cuh, answering the "what are thr_layout / val_layout" question from issue [Question] How to Understand the thr_layout and val_layout in make_tiled_copy_V_interleave_trans? #44.
  • Document the cuTe layout algebra for the FP8 V-tile smem→reg copy:
    • thr_layout = (32,4,1,1):(4,1,0,0) — 128 threads as 4 warps striding by 4 along N; each warp owns a (16,8) sub-tile of the 64-wide N.
    • val_layout = (2,2,1,4):(1,2,4,4) — 16 values per thread spanning 8 columns at 2 row-positions; the d*4 stride produces the V0/V1 interleave consumed by SM75_U16x8_LDSM_T.
    • raked_product / right_inverse mapping — how (M,N) smem coords are routed to (thread_id, value_id) so cute::copy can drive the LDSM fragment correctly.
  • Clarify that the V0/V1 look-ahead pattern comes from the strided thread traversal (4 warps split across register halves), not from val_layout strides alone.
  • Pure documentation change; no code modified.

Testing

  • No functional testing required — comments do not enter the compiled object.
  • Existing binary is bit-for-bit unaffected.

Tencent#44)

Answers the "what are thr_layout / val_layout" question from issue Tencent#44 by
documenting the cuTe layout algebra for the FP8 V-tile smem→reg copy:
  - thr_layout (32,4,1,1):(4,1,0,0): 128 threads as 4 warps striding by 4
    along N; each warp owns a (16,8) sub-tile of the 64-wide N.
  - val_layout (2,2,1,4):(1,2,4,4): 16 values per thread spanning 8
    columns at 2 row-positions; the d*4 stride produces the V0/V1 interleave
    consumed by SM75_U16x8_LDSM_T (LDSM.1matrix).
  - raked_product / right_inverse mapping: how (M,N) smem coords are routed
    to (thread_id, value_id) so cute::copy can drive the LDSM fragment.

No code changes; no rebuild required.

Co-Authored-By: Claude <noreply@anthropic.com>
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