Skip to content

[Refactor] Support auto swizzling for tma store and phaseout related layout annotations - #1509

Merged
LeiWang1999 merged 14 commits into
tile-ai:mainfrom
LeiWang1999:refactor_layout_1223
Dec 24, 2025
Merged

LeiWang1999 merged 14 commits into
tile-ai:mainfrom
LeiWang1999:refactor_layout_1223

Conversation

@LeiWang1999

@LeiWang1999 LeiWang1999 commented Dec 23, 2025 •

Copy link
Copy Markdown
Member

This pull request removes the usage of swizzled memory layouts and related annotations throughout the codebase, simplifying tensor allocation and layout management. It also makes minor improvements to variable allocation and thread binding in some kernels. These changes help streamline the code and reduce unnecessary complexity around layout specification.

Memory layout and annotation simplification:

  • Removed all calls to T.annotate_layout and the use of make_swizzled_layout for shared tensors in convolution, attention, and casting examples, as well as benchmark scripts. This affects files such as example_conv_analyze.py, example_convolution.py, example_convolution_autotune.py, example_gqa_sink_bwd_bhsd.py, example_gqa_sink_fwd_bhsd_wgmma_pipelined.py, example_mha_sink_bwd_bhsd.py, example_mha_sink_fwd_bhsd.py, example_mha_sink_fwd_bhsd_wgmma_pipelined.py, example_deepseek_mla/example_mla_decode.py, example_deepseek_mla/example_mla_decode_paged.py, and example_deepseek_mla/example_mla_decode_persistent.py. [1] [2] [3] [4] [5] [6] [7] [8] [9] [10] [11] [12] [13] [14] [15] [16] [17] [18] [19] [20]

  • Removed imports of make_swizzled_layout and tilelang.layout where they are no longer needed, further cleaning up the code. [1] [2] [3] [4] [5] [6]

Tensor allocation and scope changes:

  • Updated tensor allocation in benchmark_mamba_chunk_scan.py to remove the explicit scope="shared" and scope="shared.dyn" parameters, using simpler allocation calls instead.

Kernel and utility improvements:

  • Changed thread binding in tilelang_bitnet_158_int8xint2_prefill.py to use T.get_thread_binding(0) instead of T.thread_binding(0, threads, "threadIdx.x"), simplifying thread management.

  • Improved variable allocation in example_group_per_split_token_cast_to_fp8.py by replacing fragment allocation for row_offset with a simple scalar variable, and updated its usage accordingly.

Miscellaneous:

  • Adjusted the logic in get_bwd_configs() in example_gqa_sink_bwd_bhsd.py to always return the default configuration for SM version 90 and above, removing the exception for unsupported versions.

Summary by CodeRabbit

  • Refactor

    • Removed many explicit swizzled/layout annotations so automatic layout inference is used.
    • Simplified shared-memory allocation/scope usage across kernels and examples.
    • Safer 1D bulk/TMA paths with conservative fallbacks.
  • Bug Fixes

    • Stricter layout validation for certain reductions to prevent invalid non-linear layouts.
  • Chores

    • Added readable enum/string helpers for diagnostics.
    • Public API: linear-layout constructor now accepts a shape-based input.
    • Deleted one obsolete example.
  • Tests

    • Disabled a flaky CUDA postprocessing test.

✏️ Tip: You can customize this high-level summary in your review settings.

… sink examples and kernel analysis script for improved clarity and performance.
@github-actions

Copy link
Copy Markdown

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run pre-commit run --all-files in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀

@coderabbitai

coderabbitai Bot commented Dec 23, 2025 •

Copy link
Copy Markdown
Contributor

Warning

Rate limit exceeded

@LeiWang1999 has exceeded the limit for the number of commits that can be reviewed per hour. Please wait 7 minutes and 0 seconds before requesting another review.

⌛ How to resolve this issue?

After the wait time has elapsed, a review can be triggered using the @coderabbitai review command as a PR comment. Alternatively, push new commits to this PR.

We recommend that you space out your commits to avoid hitting the rate limit.

🚦 How do rate limits work?

CodeRabbit enforces hourly rate limits for each developer per organization.

Our paid plans have higher rate limits than the trial, open-source and free plans. In all cases, we re-allow further reviews after a brief timeout.

Please see our FAQ for further information.

📥 Commits

Reviewing files that changed from the base of the PR and between 5a0f025 and ec4a285.

📒 Files selected for processing (2)
  • examples/dequantize_gemm/regression_example_dequantize_gemm.py
  • examples/dequantize_gemm/test_example_dequantize_gemm.py

Note

Other AI code review bot(s) detected

CodeRabbit has detected other AI code review bot(s) in this pull request and will avoid duplicating their findings in the review comments. This may lead to a less comprehensive review.

📝 Walkthrough

Walkthrough

Removed many explicit TileLang swizzled-layout annotations and several shared-allocation scope arguments across examples and benchmarks; introduced conservative 1D BulkLoad/BulkStore handling with 1D‑TMA guards and fallbacks in copy lowering; added enum-to-string debug helpers and replaced the two-arg linear-layout API with a shape-based linear layout.

Changes

Cohort / File(s) Summary
Examples — attention / flash / mha / gqa / deepseek / linear
examples/attention_sink/..., examples/flash_attention/..., examples/mha_*, examples/gqa_*, examples/deepseek_*, examples/linear_attention/*
Removed numerous T.annotate_layout swizzled-layout mappings for Q/K/V/O/S and related shared fragments; removed imports of swizzle helpers where present.
Examples — convolution / analyze / tests
examples/convolution/..., examples/analyze/example_conv_analyze.py, testing/python/jit/test_tilelang_jit_*.py
Deleted swizzled-layout annotations for out_shared, data_shared, kernel_shared (tests updated similarly); removed make_swizzled_layout imports and annotation blocks.
Examples — dequantize / grouped GEMM / deleted file
examples/dequantize_gemm/...
Removed swizzled-layout annotations for various A/B/C/Ct buffers; deleted example_dequant_gemm_bf16_mxfp4_hopper_tma.py entirely; removed some Ct_shared/Ct usages and related annotations.
Examples — GDN / DSA sparse / MLA / chunk kernels
examples/gdn/..., examples/dsa_sparse_finetune/..., examples/deepseek_nsa/..., examples/linear_attention/example_mamba_chunk_*.py
Dropped multiple shared-buffer swizzle annotations; in mamba chunk examples removed explicit scope args from T.alloc_shared(...) calls and removed acc_o_shared from T.annotate_layout mappings.
Benchmarks — matmul / mamba
benchmark/matmul/..., benchmark/matmul_fp8/..., benchmark/mamba2/benchmark_mamba_chunk_scan.py
Removed top-level tilelang imports in some files and removed explicit swizzled-layout annotations for C_shared; removed allocation scope args and dropped acc_o_shared from layout maps in mamba benchmark.
Examples — casting / thread binding
examples/cast/*, examples/bitnet-1.58b/...
Replaced per-fragment row_offset with scalar row_offset in a cast example; switched thread-binding API from T.thread_binding(...) to T.get_thread_binding(...).
Examples — flash_decoding / combines / local buffers
examples/flash_decoding/...
Removed per-thread layout annotate blocks and some intermediate shared allocations (e.g., po_shared), replacing certain copies to use local buffers instead of removed shared buffers.
Core — copy lowering & 1D TMA/layout inference
src/op/copy.cc
Extended Copy lowering and InferLayout to handle 1D BulkLoad/BulkStore (BulkLoad1D/BulkStore1D) with is_tma_1d and shared_range guards, early short-circuit for single-dim cases, conservative 1D swizzle computation, and fallbacks to non-TMA lowering.
Core headers — debug helpers
src/op/copy.h, src/op/gemm.h, src/op/operator.h
Added inline enum-to-string helpers: CopyInstToString, GemmWarpPolicyTypeToString, GemmInstToString, and InferLevelToString for logging/debugging.
Layout API & bindings
src/layout/gemm_layouts.cc, src/layout/layout.cc, src/layout/layout.h, tilelang/layout/swizzle.py
Replaced the two-arg linear-layout helper with makeLinearLayout(Array<PrimExpr> shape); updated tl.make_linear_layout binding to accept a shape array; Python wrapper make_linear_layout now accepts a buffer-like object and derives shape. Call sites updated.
Reduce layout inference
src/op/reduce.cc
Tightened CumSumOpNode::InferLayout to infer/validate linear layouts for shared buffers in strict mode and error on non-linear shared layouts.
Tests / small changes
testing/python/jit/test_tilelang_jit_callback.py, testing/python/jit/test_tilelang_jit_nvrtc.py
Added pytest import and skipped a CUDA postproc callback test; removed swizzle-layout annotate blocks from NVRTC/FFI tests.

Sequence Diagram(s)

(silently omitted)

Estimated code review effort

🎯 4 (Complex) | ⏱️ ~45 minutes

Possibly related PRs

Poem

"I nibble code with eager paws,
Removed some swizzles, fixed the laws,
Shared buffers now breathe free,
1D paths hop cautiously,
A tiny rabbit's celebratory cheer 🐇"

Pre-merge checks and finishing touches

❌ Failed checks (1 warning)
Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 49.33% which is insufficient. The required threshold is 80.00%. You can run @coderabbitai generate docstrings to improve docstring coverage.
✅ Passed checks (2 passed)
Check name Status Explanation
Description Check ✅ Passed Check skipped - CodeRabbit’s high-level summary is enabled.
Title check ✅ Passed The PR title '[Refactor] Support auto swizzling for tma store and phaseout related layout annotations' directly summarizes the main change: removing explicit layout annotations and supporting automatic swizzling instead.

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands and usage tips.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 2

🧹 Nitpick comments (2)
examples/attention_sink/example_gqa_sink_bwd_bhsd.py (1)

11-17: LGTM: Simplified SM version handling.

The change removes the exception for unsupported SM versions and uses a default configuration for SM >= 90, which aligns with the PR objectives. The larger config (block_M=128, num_stages=2, threads=256) is appropriate for newer architectures with more resources.

Consider adding a brief comment explaining why SM 80 uses different configuration parameters:

Optional: Add explanatory comment
 def get_bwd_configs():
+    # SM 80 (A100) uses smaller blocks due to resource constraints
+    # SM 90+ (H100, etc.) uses larger blocks for better performance
     sm_major, sm_minor = torch.cuda.get_device_capability()
     sm_version = sm_major * 10 + sm_minor
     if sm_version == 80:
         return 64, 32, 1, 128
     else:
         return 128, 32, 2, 256
examples/linear_attention/example_mamba_chunk_scan.py (1)

122-122: Consider removing redundant scope parameter for consistency.

Line 122 still has an explicit scope="shared.dyn" parameter, while similar allocations at lines 115, 117, 120, 123, and 128 had their scope parameters removed in this refactoring. Since "shared.dyn" is the default for T.alloc_shared, this explicit parameter is redundant.

🔎 Proposed change for consistency
-            x_shared = T.alloc_shared((block_K, block_N), dtype, scope="shared.dyn")
+            x_shared = T.alloc_shared((block_K, block_N), dtype)
📜 Review details

Configuration used: defaults

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between e79bbcc and bd317fa.

📒 Files selected for processing (59)
  • benchmark/mamba2/benchmark_mamba_chunk_scan.py
  • benchmark/matmul/benchmark_matmul.py
  • benchmark/matmul_fp8/benchmark_matmul.py
  • examples/analyze/example_conv_analyze.py
  • examples/attention_sink/example_gqa_sink_bwd_bhsd.py
  • examples/attention_sink/example_gqa_sink_fwd_bhsd_wgmma_pipelined.py
  • examples/attention_sink/example_mha_sink_bwd_bhsd.py
  • examples/attention_sink/example_mha_sink_fwd_bhsd.py
  • examples/attention_sink/example_mha_sink_fwd_bhsd_wgmma_pipelined.py
  • examples/bitnet-1.58b/kernel_benchmark/tilelang_bitnet_158_int8xint2_prefill.py
  • examples/cast/example_group_per_split_token_cast_to_fp8.py
  • examples/cast/example_per_token_cast_to_fp8.py
  • examples/convolution/example_convolution.py
  • examples/convolution/example_convolution_autotune.py
  • examples/deepseek_mla/example_mla_decode.py
  • examples/deepseek_mla/example_mla_decode_paged.py
  • examples/deepseek_mla/example_mla_decode_persistent.py
  • examples/deepseek_mla/experimental/example_mla_decode_kv_fp8.py
  • examples/deepseek_nsa/benchmark/benchmark_nsa_fwd.py
  • examples/deepseek_nsa/example_tilelang_nsa_bwd.py
  • examples/deepseek_v32/sparse_mla_bwd.py
  • examples/dequantize_gemm/example_dequant_gemm_bf16_fp4_hopper.py
  • examples/dequantize_gemm/example_dequant_gemm_bf16_mxfp4_hopper.py
  • examples/dequantize_gemm/example_dequant_gemm_bf16_mxfp4_hopper_tma.py
  • examples/dequantize_gemm/example_dequant_gemm_fp4_hopper.py
  • examples/dequantize_gemm/example_dequant_gemm_w4a8.py
  • examples/dequantize_gemm/example_dequant_groupedgemm_bf16_mxfp4_hopper.py
  • examples/dsa_sparse_finetune/sparse_mla_bwd.py
  • examples/flash_attention/README.md
  • examples/flash_attention/example_gqa_bwd.py
  • examples/flash_attention/example_gqa_bwd_tma_reduce.py
  • examples/flash_attention/example_gqa_bwd_tma_reduce_varlen.py
  • examples/flash_attention/example_gqa_bwd_wgmma_pipelined.py
  • examples/flash_attention/example_gqa_fwd_varlen.py
  • examples/flash_attention/example_mha_bwd_bhsd.py
  • examples/flash_attention/example_mha_bwd_bshd_wgmma_pipelined.py
  • examples/flash_attention/example_mha_fwd_varlen.py
  • examples/flash_decoding/example_gqa_decode.py
  • examples/flash_decoding/example_gqa_decode_varlen_logits.py
  • examples/flash_decoding/example_mha_inference.py
  • examples/gdn/example_chunk_delta_bwd.py
  • examples/gdn/example_chunk_delta_h.py
  • examples/gdn/example_chunk_o.py
  • examples/gdn/example_chunk_o_bwd.py
  • examples/gdn/example_chunk_scaled_dot_kkt.py
  • examples/gdn/example_wy_fast.py
  • examples/linear_attention/example_linear_attn_bwd.py
  • examples/linear_attention/example_linear_attn_fwd.py
  • examples/linear_attention/example_mamba_chunk_scan.py
  • examples/linear_attention/example_mamba_chunk_state.py
  • examples/minference/example_vertical_slash_sparse_attn.py
  • examples/warp_specialize/example_warp_specialize_flashmla.py
  • src/op/copy.cc
  • src/op/copy.h
  • src/op/gemm.h
  • src/op/operator.h
  • testing/python/jit/test_tilelang_jit_nvrtc.py
  • testing/python/jit/test_tilelang_jit_tvm_ffi.py
  • testing/python/language/test_tilelang_language_vectorized_cast.py
💤 Files with no reviewable changes (45)
  • examples/dequantize_gemm/example_dequant_gemm_w4a8.py
  • examples/minference/example_vertical_slash_sparse_attn.py
  • examples/gdn/example_chunk_delta_h.py
  • examples/deepseek_mla/example_mla_decode_persistent.py
  • examples/gdn/example_chunk_o_bwd.py
  • examples/flash_decoding/example_gqa_decode_varlen_logits.py
  • examples/convolution/example_convolution.py
  • examples/flash_decoding/example_gqa_decode.py
  • examples/warp_specialize/example_warp_specialize_flashmla.py
  • testing/python/jit/test_tilelang_jit_nvrtc.py
  • examples/linear_attention/example_linear_attn_fwd.py
  • examples/dequantize_gemm/example_dequant_gemm_bf16_fp4_hopper.py
  • examples/dsa_sparse_finetune/sparse_mla_bwd.py
  • examples/deepseek_nsa/benchmark/benchmark_nsa_fwd.py
  • examples/dequantize_gemm/example_dequant_gemm_bf16_mxfp4_hopper.py
  • benchmark/matmul_fp8/benchmark_matmul.py
  • examples/dequantize_gemm/example_dequant_gemm_bf16_mxfp4_hopper_tma.py
  • examples/flash_attention/example_mha_bwd_bshd_wgmma_pipelined.py
  • examples/flash_attention/example_gqa_fwd_varlen.py
  • examples/gdn/example_chunk_o.py
  • examples/deepseek_v32/sparse_mla_bwd.py
  • examples/convolution/example_convolution_autotune.py
  • examples/dequantize_gemm/example_dequant_groupedgemm_bf16_mxfp4_hopper.py
  • examples/deepseek_mla/example_mla_decode.py
  • examples/flash_attention/example_mha_fwd_varlen.py
  • testing/python/jit/test_tilelang_jit_tvm_ffi.py
  • examples/attention_sink/example_mha_sink_fwd_bhsd_wgmma_pipelined.py
  • examples/flash_attention/example_mha_bwd_bhsd.py
  • examples/cast/example_per_token_cast_to_fp8.py
  • examples/deepseek_mla/example_mla_decode_paged.py
  • examples/attention_sink/example_mha_sink_fwd_bhsd.py
  • examples/attention_sink/example_mha_sink_bwd_bhsd.py
  • examples/flash_attention/example_gqa_bwd.py
  • examples/flash_attention/example_gqa_bwd_tma_reduce.py
  • examples/gdn/example_wy_fast.py
  • examples/dequantize_gemm/example_dequant_gemm_fp4_hopper.py
  • examples/linear_attention/example_linear_attn_bwd.py
  • examples/gdn/example_chunk_scaled_dot_kkt.py
  • examples/flash_attention/example_gqa_bwd_wgmma_pipelined.py
  • examples/deepseek_nsa/example_tilelang_nsa_bwd.py
  • examples/flash_attention/example_gqa_bwd_tma_reduce_varlen.py
  • examples/deepseek_mla/experimental/example_mla_decode_kv_fp8.py
  • examples/analyze/example_conv_analyze.py
  • examples/gdn/example_chunk_delta_bwd.py
  • benchmark/matmul/benchmark_matmul.py
🧰 Additional context used
🧠 Learnings (2)
📓 Common learnings
Learnt from: botbw
Repo: tile-ai/tilelang PR: 691
File: src/tl_templates/cuda/gemm_sp_sm80.h:81-85
Timestamp: 2025-09-15T10:51:06.985Z
Learning: In CUTLASS tensor operation layouts, crosswise constants should be computed using sizeof(T) (bytes), not cutlass::sizeof_bits<T>::value (bits). This is the established pattern in the official CUTLASS codebase, as seen in default_mma_core_sparse_sm80.h.
Learnt from: botbw
Repo: tile-ai/tilelang PR: 691
File: src/tl_templates/cuda/gemm_sp_sm80.h:81-85
Timestamp: 2025-09-15T10:51:06.985Z
Learning: In CUTLASS tensor operation layouts, crosswise constants should be computed using sizeof(T) (bytes), not cutlass::sizeof_bits<T>::value (bits). However, the layout template parameter should use sizeof_bits<T>::value (bits). This is the established pattern in the official CUTLASS codebase, as seen in default_mma_core_sparse_sm80.h where Crosswise uses sizeof(ElementA) but the layout template uses sizeof_bits<ElementA>::value.
📚 Learning: 2025-11-14T07:56:11.098Z
Learnt from: lucifer1004
Repo: tile-ai/tilelang PR: 1256
File: testing/python/jit/test_tilelang_jit_gemm_nvrtc.py:55-115
Timestamp: 2025-11-14T07:56:11.098Z
Learning: In `testing/python/jit/test_tilelang_jit_gemm_nvrtc.py`, the global function `tilelang_callback_cuda_postproc` registered via `tvm.register_global_func(..., override=True)` is intentionally not restored after the test completes, as the persistent behavior is expected.

Applied to files:

  • testing/python/language/test_tilelang_language_vectorized_cast.py
🧬 Code graph analysis (7)
testing/python/language/test_tilelang_language_vectorized_cast.py (1)
tilelang/language/v2/dtypes.py (2)
  • float16 (294-294)
  • float4_e2m1fn (385-385)
examples/linear_attention/example_mamba_chunk_state.py (2)
tilelang/language/annotations.py (1)
  • annotate_layout (26-37)
tilelang/layout/swizzle.py (1)
  • make_swizzled_layout (62-71)
src/op/gemm.h (1)
tilelang/tileop/gemm/__init__.py (1)
  • GemmInst (33-52)
examples/attention_sink/example_gqa_sink_fwd_bhsd_wgmma_pipelined.py (1)
tilelang/env.py (1)
  • disable_cache (272-273)
benchmark/mamba2/benchmark_mamba_chunk_scan.py (1)
tilelang/language/allocate.py (2)
  • alloc_shared (40-55)
  • alloc_fragment (72-85)
examples/cast/example_group_per_split_token_cast_to_fp8.py (1)
tilelang/language/v2/dtypes.py (1)
  • int32 (240-240)
examples/linear_attention/example_mamba_chunk_scan.py (1)
tilelang/language/allocate.py (2)
  • alloc_shared (40-55)
  • alloc_fragment (72-85)
⏰ Context from checks skipped due to timeout of 90000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (3)
  • GitHub Check: Test for Python 3.12 with Nightly-ROCm-7.1 (on self-hosted-amd)
  • GitHub Check: Test for Python 3.12 with Metal (on macos-latest)
  • GitHub Check: Test for Python 3.12 with CUDA-12.8 (on self-hosted-nvidia)
🔇 Additional comments (15)
examples/flash_attention/README.md (1)

1-109: Documentation update aligns well with the refactoring objectives.

The removal of explicit T.annotate_layout calls for shared buffers (Q_shared, K_shared, V_shared) is consistent with the PR's goal of simplifying layout management by supporting auto-swizzling. The documented example remains accurate and functional—the buffers are still properly allocated in shared memory, just without manual layout annotations. The code flow and kernel logic remain unchanged and well-explained.

examples/cast/example_group_per_split_token_cast_to_fp8.py (3)

34-34: LGTM: Cleaner scalar allocation.

The change from T.alloc_fragment((1,), T.int32) to T.alloc_var(dtype=T.int32) appropriately simplifies the allocation for a scalar offset value, improving code readability.


36-38: LGTM: Consistent scalar usage.

The initialization and accumulation logic correctly use the scalar variable. The loop properly computes the cumulative offset by summing batch sizes from index 0 to bg-1.


41-41: LGTM: Simplified array indexing.

The slice expression correctly uses the scalar offset, removing the unnecessary [0] indexing and improving readability.

examples/linear_attention/example_mamba_chunk_scan.py (1)

115-115: LGTM - Scope parameter removal aligns with PR objectives.

The removal of explicit scope parameters from these T.alloc_shared calls is correct. They now default to scope="shared.dyn" as defined in the alloc_shared function signature, which is the intended behavior for this refactoring.

Also applies to: 117-117, 120-120, 123-123, 128-128

testing/python/language/test_tilelang_language_vectorized_cast.py (1)

130-131: Revert to standard test harness pattern.

The test harness tilelang.testing.main() is commented out and replaced with a hard-coded call to a single test function. This deviates from the standard pattern used by all other test files in the repository, which call tilelang.testing.main() to discover and run tests via pytest.

Restore line 130 to tilelang.testing.main() and remove the hard-coded test call on line 131.

Likely an incorrect or invalid review comment.

src/op/operator.h (1)

35-47: LGTM!

Clean utility function for debugging. All current InferLevel enum cases are covered, and the default case handles potential future enum additions gracefully.

src/op/copy.h (1)

31-55: LGTM!

All CopyInst enum values are covered with appropriate string representations. Consistent pattern with other *ToString helpers added in this PR.

src/op/gemm.h (1)

25-58: LGTM!

Both helper functions correctly cover all enum cases for GemmWarpPolicyType and GemmInst. The string representations align with the Python enum's __repr__ behavior.

src/op/copy.cc (1)

548-573: LGTM on 1D TMA condition handling.

The expanded condition correctly handles the new kBulkLoad1D and kBulkStore1D variants. The early return for single-dimension 1D TMA store is appropriate since swizzling cannot be applied to truly 1D layouts.

examples/flash_decoding/example_mha_inference.py (1)

183-184: LGTM!

Simplifying the data path by copying directly from Output_partial to po_local eliminates the intermediate shared buffer (po_shared). This aligns with the PR's objective of removing explicit layout annotations and relying on auto-inferred layouts.

examples/attention_sink/example_gqa_sink_fwd_bhsd_wgmma_pipelined.py (1)

146-149: The simplified shared buffer allocations are correct and align with the PR objective of supporting automatic layout inference for TMA operations. The compiler's layout inference assigns appropriate layouts to each buffer based on context, and swizzling strategies are automatically handled by TileLang, resulting in no visible differences in the code. This approach reduces verbosity while maintaining performance through compiler-driven optimization.

benchmark/mamba2/benchmark_mamba_chunk_scan.py (2)

251-256: Clarify whether layout annotations should be fully removed.

The PR objectives indicate removing explicit swizzled layouts and layout annotations across the codebase. However, this file still retains T.annotate_layout with make_swizzled_layout calls for cb_shared and x_residual_shared.

Is this intentional? Should these annotations also be removed as part of the auto swizzling support, or do these specific tensors require explicit layout specifications?


228-241: Code state is correct: T.alloc_shared calls properly default to dynamic shared memory.

All T.alloc_shared calls in lines 228-241 correctly omit explicit scope parameters, relying on the default scope="shared.dyn" which supports auto swizzling. Since the kernel uses float16 and float32 dtypes (not bool), the dynamic scope applies uniformly across all shared memory allocations.

Layout annotations (lines 251-256) are appropriately applied only to cb_shared and x_residual_shared, consistent with the swizzling optimization strategy.

examples/linear_attention/example_mamba_chunk_state.py (1)

92-92: Clarify the rationale for asymmetric layout annotation removal.

Only x_shared retains its swizzled layout annotation (line 92), while acc_o_shared has no annotation. Since x_shared is used in load operations (lines 104, 112) and acc_o_shared in store operations (lines 125-129), please confirm:

  1. Is this asymmetry intentional—does auto-swizzling now cover store paths to acc_o_shared while loads still require manual annotation for x_shared?
  2. Has removing the acc_o_shared annotation been performance-validated against the previous version?
  3. Are both buffers covered by the auto-optimization mentioned in the PR title, or is there a plan to eventually remove the x_shared annotation once auto-swizzling fully covers load operations?

Comment on lines +13 to +14
tilelang.disable_cache()

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

⚠️ Potential issue | 🟠 Major

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
# Description: Check for cache disable patterns across attention examples

# Search for disable_cache calls in example files
echo "=== Searching for disable_cache in examples ==="
rg -n "disable_cache" examples/

echo -e "\n=== Searching for cache-related calls in attention examples ==="
fd "example.*attention.*\.py$" examples/ 2>/dev/null | head -20

echo -e "\n=== Check for cache-related patterns in all example files ==="
rg -l "cache" examples/ --type py | head -10

echo -e "\n=== Check context around the file in question ==="
if [ -f "examples/attention_sink/example_gqa_sink_fwd_bhsd_wgmma_pipelined.py" ]; then
    head -20 examples/attention_sink/example_gqa_sink_fwd_bhsd_wgmma_pipelined.py
fi

Repository: tile-ai/tilelang

Length of output: 2264


🏁 Script executed:

#!/bin/bash
# Check full file structure and autotune usage
echo "=== Full file content ==="
cat examples/attention_sink/example_gqa_sink_fwd_bhsd_wgmma_pipelined.py

echo -e "\n=== Check if test imports this file ==="
grep -r "example_gqa_sink_fwd_bhsd_wgmma_pipelined" examples/ tests/

echo -e "\n=== Check if this file has __main__ guard ==="
grep -n "__main__" examples/attention_sink/example_gqa_sink_fwd_bhsd_wgmma_pipelined.py

echo -e "\n=== Compare with flash_attention example (which also has disable_cache) ==="
head -25 examples/flash_attention/example_gqa_bwd_tma_reduce.py

Repository: tile-ai/tilelang

Length of output: 16540


Move cache disable to controlled execution scope.

The cache is disabled at module import time, creating a global side effect for all code that imports this module—including test files. Since the file has a __name__ == "__main__" guard, move tilelang.disable_cache() into that block or into the benchmark/tune functions that require it, rather than executing it unconditionally on import.

🤖 Prompt for AI Agents
In examples/attention_sink/example_gqa_sink_fwd_bhsd_wgmma_pipelined.py around
lines 13-14, tilelang.disable_cache() is called at import time causing a global
side effect; move this call into a controlled execution scope such as inside the
existing if __name__ == "__main__" block or into the specific benchmark/tune
functions that need it so importing the module does not change global cache
state. Ensure the call executes only when running the script or when the
specific functions are invoked, not during module import.

Comment thread src/op/copy.cc Resolved
@LeiWang1999

Copy link
Copy Markdown
Member Author

@regression-perf

@github-actions

Copy link
Copy Markdown

Performance Benchmark Report

Triggered by: @LeiWang1999
Workflow run: https://github.com/tile-ai/tilelang/actions/runs/20456794913

Results

File Original Latency Current Latency Speedup
example_dequant_gemm_bf16_mxfp4_hopper_tma 0.012768 0.019424 0.657331
example_warp_specialize_gemm_barrierpipe_stage2 0.038849 0.040097 0.968875
example_warp_specialize_gemm_softpipe_stage2 0.037792 0.038849 0.972792
example_dequant_gemm_w4a8 0.005312 0.005344 0.994012
example_mha_inference 0.0765904 0.0768307 0.996872
example_gemm_schedule 0.0325711 0.0326404 0.997874
example_elementwise_add 0.0224843 0.0225203 0.998401
example_gemm_intrinsics 0.035456 0.035488 0.999098
sparse_mla_bwd 0.381521 0.381619 0.999745
example_fusedmoe_tilelang 0.147391 0.147401 0.999936
example_gemm 0.022784 0.022785 0.999956
example_linear_attn_fwd 0.0350014 0.0350019 0.999984
example_dequant_gemm_fp4_hopper 0.012352 0.012352 1
topk_selector 0.0545977 0.0545934 1.00008
example_tilelang_gemm_splitk_vectorize_atomicadd 0.0497302 0.0497194 1.00022
example_linear_attn_bwd 0.144296 0.144245 1.00036
tilelang_example_sparse_tensorcore 0.0152087 0.0152006 1.00053
example_gemv 0.0670788 0.0670253 1.0008
example_dequant_gemv_fp16xint4 0.00379653 0.00379307 1.00091
example_dequant_groupedgemm_bf16_mxfp4_hopper 0.0172939 0.0172701 1.00138
example_tilelang_gemm_splitk 0.0497035 0.0496264 1.00155
example_convolution_autotune 0.999213 0.99516 1.00407
example_dequant_gemm_bf16_fp4_hopper 0.012064 0.012 1.00533
example_topk 0.011201 0.011136 1.00584
sparse_mla_fwd 0.42549 0.422784 1.0064
example_per_token_cast_to_fp8 0.00771038 0.00765993 1.00659
example_tilelang_nsa_fwd 0.00715595 0.00710314 1.00743
example_tilelang_nsa_decode 0.00683537 0.00677491 1.00892
example_mha_sink_fwd_bhsd_wgmma_pipelined 0.0157547 0.0156095 1.0093
example_mha_sink_fwd_bhsd_wgmma_pipelined_sliding_window 0.0158415 0.0156818 1.01018
example_gqa_decode 0.049409 0.048833 1.0118
example_tilelang_sparse_gqa_decode_varlen_mask 0.0239743 0.0236281 1.01465
example_tilelang_sparse_gqa_decode_varlen_indice 0.0175418 0.0172825 1.015
example_vertical_slash_sparse_attn 0.244708 0.240972 1.0155
fp8_lighting_indexer 0.0368113 0.0362162 1.01643
example_mha_sink_fwd_bhsd_sliding_window 0.0131539 0.0129149 1.0185
example_group_per_split_token_cast_to_fp8 0.0107225 0.0105259 1.01867
example_mha_sink_fwd_bhsd 0.0128814 0.0125997 1.02236
example_gqa_sink_bwd_bhsd_sliding_window 0.0256605 0.0250323 1.02509
example_convolution 1.38381 1.34758 1.02688
sparse_mla_fwd_pipelined 0.128764 0.125293 1.0277
example_mla_decode 0.480806 0.467497 1.02847
example_warp_specialize_gemm_copy_0_gemm_1 0.039617 0.038496 1.02912
example_dequant_gemm_bf16_mxfp4_hopper 0.012897 0.01248 1.03341
example_gqa_sink_fwd_bhsd_wgmma_pipelined 0.0150777 0.014578 1.03428
example_gqa_sink_fwd_bhsd_wgmma_pipelined_sliding_window 0.0152431 0.0147273 1.03502
example_mha_sink_bwd_bhsd_sliding_window 0.0344316 0.0332493 1.03556
example_mha_sink_bwd_bhsd 0.0582127 0.0559355 1.04071
block_sparse_attn_tilelang 0.0108088 0.0103646 1.04286
example_gqa_sink_bwd_bhsd 0.0429825 0.0410798 1.04632
example_warp_specialize_gemm_copy_1_gemm_0 0.039296 0.036961 1.06317
example_tilelang_block_sparse_attn 0.00847558 0.0079187 1.07032
example_blocksparse_gemm 0.0245965 0.0228276 1.07749
example_gemm_autotune 0.025665 0.02272 1.12962

Artifacts

  • regression_result.png (speedup plot) is attached as a workflow artifact. Download it from the workflow run page above.

This change introduces a layout annotation for the B_shared tensor in the example_dequant_gemm_bf16_mxfp4_hopper.py file, enhancing the memory layout optimization for better performance during matrix multiplication.
@LeiWang1999

Copy link
Copy Markdown
Member Author

@codex review

@chatgpt-codex-connector chatgpt-codex-connector Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💡 Codex Review

Here are some automated review suggestions for this pull request.

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

  • Open a pull request for review
  • Mark a draft as ready
  • Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

Comment thread src/op/copy.cc Resolved
…dimension

- Renamed `makeGemmLayoutLinear` to `makeLinearLayout` and updated its implementation to handle arbitrary dimensions.
- Updated related function calls in `gemm_layouts.cc` and `layout.cc` to use the new layout function.
- Enhanced layout inference in `CumSumOpNode` to enforce linear layout for shared buffers in strict mode.
… arbitrary dimensions

- Updated the function signature to take a `Buffer`, `BufferLoad`, or `BufferRegion` directly.
- Simplified the implementation by removing argument checks and directly obtaining the shape from the buffer info.
- Enhanced the documentation to clarify the function's purpose and usage.
@LeiWang1999

Copy link
Copy Markdown
Member Author

@regression-perf

@github-actions

Copy link
Copy Markdown

Performance Benchmark Report

Triggered by: @LeiWang1999
Workflow run: https://github.com/tile-ai/tilelang/actions/runs/20467645826

Results

File Original Latency Current Latency Speedup
example_dequant_gemm_bf16_mxfp4_hopper_tma 0.01248 0.01952 0.639344
example_warp_specialize_gemm_softpipe_stage2 0.037505 0.038977 0.962234
example_dequant_gemm_w4a8 0.00528 0.005376 0.982143
example_topk 0.011104 0.011296 0.983003
example_warp_specialize_gemm_barrierpipe_stage2 0.039169 0.039712 0.986327
sparse_mla_fwd_pipelined 0.139802 0.141506 0.98796
example_dequant_gemm_fp4_hopper 0.012353 0.012449 0.992289
example_mha_inference 0.0764309 0.0766066 0.997706
example_elementwise_add 0.0224589 0.0225052 0.99794
example_tilelang_gemm_splitk_vectorize_atomicadd 0.0496507 0.049704 0.998927
example_linear_attn_bwd 0.144238 0.144334 0.999334
tilelang_example_sparse_tensorcore 0.0152025 0.0152054 0.999806
sparse_mla_bwd 0.381538 0.38158 0.99989
example_gemm_intrinsics 0.035488 0.035489 0.999972
example_linear_attn_fwd 0.0349946 0.0349944 1
example_vertical_slash_sparse_attn 0.240991 0.240986 1.00002
example_fusedmoe_tilelang 0.147408 0.147388 1.00013
example_dequant_gemv_fp16xint4 0.00379808 0.00379684 1.00033
example_gemm_schedule 0.0326485 0.0326326 1.00049
topk_selector 0.0546212 0.0545734 1.00088
example_tilelang_gemm_splitk 0.0497121 0.0496106 1.00205
example_gemv 0.0671643 0.0670148 1.00223
example_dequant_groupedgemm_bf16_mxfp4_hopper 0.0172408 0.0172002 1.00236
example_dequant_gemm_bf16_fp4_hopper 0.01216 0.012096 1.00529
example_gemm 0.022913 0.022785 1.00562
sparse_mla_fwd 0.755362 0.751088 1.00569
example_convolution_autotune 1.00005 0.994365 1.00572
example_per_token_cast_to_fp8 0.00770583 0.00766177 1.00575
example_mha_sink_bwd_bhsd_sliding_window 0.0335704 0.0333276 1.00729
example_tilelang_nsa_decode 0.00683994 0.00678053 1.00876
example_mha_sink_fwd_bhsd_sliding_window 0.013027 0.0129078 1.00923
example_tilelang_nsa_fwd 0.00717465 0.00710554 1.00973
example_gqa_decode 0.048706 0.048226 1.00995
example_mha_sink_fwd_bhsd_wgmma_pipelined_sliding_window 0.0158221 0.0156651 1.01002
example_mha_sink_fwd_bhsd 0.0127034 0.0125749 1.01022
example_mha_sink_fwd_bhsd_wgmma_pipelined 0.0157593 0.0155958 1.01048
example_tilelang_sparse_gqa_decode_varlen_mask 0.023907 0.0236272 1.01184
example_convolution 1.36653 1.34906 1.01295
example_mla_decode 0.474058 0.467498 1.01403
example_gqa_sink_fwd_bhsd_wgmma_pipelined_sliding_window 0.0149489 0.0147278 1.01501
fp8_lighting_indexer 0.0367794 0.0362257 1.01528
example_tilelang_sparse_gqa_decode_varlen_indice 0.0175518 0.0172871 1.01532
example_gqa_sink_fwd_bhsd_wgmma_pipelined 0.0148374 0.0145945 1.01665
example_group_per_split_token_cast_to_fp8 0.010729 0.0105318 1.01872
example_gqa_sink_bwd_bhsd_sliding_window 0.0255563 0.0250812 1.01894
example_gqa_sink_bwd_bhsd 0.0423058 0.0412786 1.02488
example_mha_sink_bwd_bhsd 0.0575076 0.0556909 1.03262
example_warp_specialize_gemm_copy_1_gemm_0 0.039393 0.038081 1.03445
example_dequant_gemm_bf16_mxfp4_hopper 0.012864 0.012416 1.03608
example_warp_specialize_gemm_copy_0_gemm_1 0.038881 0.037505 1.03669
block_sparse_attn_tilelang 0.0108052 0.0103622 1.04275
example_tilelang_block_sparse_attn 0.00846909 0.00790963 1.07073
example_blocksparse_gemm 0.0246268 0.022862 1.07719
example_gemm_autotune 0.025632 0.022624 1.13296

Artifacts

  • regression_result.png (speedup plot) is attached as a workflow artifact. Download it from the workflow run page above.

@LeiWang1999

Copy link
Copy Markdown
Member Author

@codex review

@chatgpt-codex-connector chatgpt-codex-connector Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💡 Codex Review

Here are some automated review suggestions for this pull request.

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

  • Open a pull request for review
  • Mark a draft as ready
  • Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

Comment on lines -1 to -5
import tilelang
import tilelang.language as T
from tilelang import tvm as tvm
from tvm import DataType
from tvm import tir

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P1 Badge Fix missing TMA example module import

Deleting this module means examples/dequantize_gemm/test_example_dequantize_gemm.py and examples/dequantize_gemm/regression_example_dequantize_gemm.py still import example_dequant_gemm_bf16_mxfp4_hopper_tma, so running those tests/regressions now raises ModuleNotFoundError during import (before any CUDA checks). If the example is intentionally removed, the corresponding imports and test functions should be removed/updated; otherwise keep a stub or restore the module.

Useful? React with 👍 / 👎.

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