---
id: pr-tilelang-2002
repo: tile-ai/tilelang
pr: 2002
title: '[Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS'
author: LeiWang1999
date: '2026-04-07'
url: https://github.com/tile-ai/tilelang/pull/2002
source_category: upstream-code
architectures:
- sm100
tags:
- mla
- mxfp4
techniques:
- pipeline-stages
hardware_features:
- fp4
- mxfp4
kernel_types:
- mla
languages: []
captured_at: '2026-05-16'
status: merged
merge_sha: 3ee0988c86536b5867a84543fc72f1271a75f013
inclusion_reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.
changed_paths:
- .gitignore
- examples/bitnet-1.58b/tokenization_bitnet.py
- examples/deepseek_mhc/example_mhc_bwd.py
- examples/deepseek_mhc/example_mhc_post.py
- examples/deepseek_mhc/example_mhc_pre.py
- examples/deepseek_mla/example_mla_decode.py
- examples/deepseek_nsa/benchmark/benchmark_nsa_fwd.py
- examples/deepseek_nsa/example_tilelang_nsa_bwd.py
- examples/deepseek_nsa/example_tilelang_nsa_decode.py
- examples/deepseek_nsa/example_tilelang_nsa_fwd.py
- examples/deepseek_nsa/example_tilelang_nsa_fwd_varlen.py
- examples/deepseek_v32/inference/convert.py
- examples/deepseek_v32/inference/kernel.py
- examples/deepseek_v32/sparse_mla_bwd.py
- examples/deepseek_v32/sparse_mla_fwd.py
- examples/dequantize_gemm/example_dequant_gemm_bf16_fp4_hopper.py
- examples/dsa_sparse_finetune/indexer_bwd.py
- examples/dsa_sparse_finetune/indexer_topk_reducesum.py
- examples/dsa_sparse_finetune/sparse_mla_bwd.py
- examples/dsa_sparse_finetune/sparse_mla_fwd.py
- examples/dsa_sparse_finetune/sparse_mla_topk_reducesum.py
- examples/flash_attention_sm100/gqa_bwd_bshd.py
- examples/flash_attention_sm100/gqa_fwd_bshd.py
- examples/flash_attention_sm100/mha_bwd_bshd.py
- examples/flash_attention_sm100/mha_fwd_bshd.py
- examples/flash_decoding/example_gqa_decode.py
- examples/fusedmoe/example_fusedmoe_tilelang.py
- examples/gdn/example_chunk_o_bwd.py
- examples/gdn/example_cumsum.py
- examples/gdn/example_wy_fast_bwd_split.py
- examples/gdn/test_example_gdn_compilation.py
- examples/gemm_fp8/example_tilelang_gemm_fp8_intrinsic.py
- examples/gemm_fp8/example_tilelang_gemm_fp8_sm100.py
- examples/gemm_sm100/README.md
- examples/gemm_sm100/gemm_mma.py
- examples/gemm_sm100/gemm_tcgen5mma.py
- examples/gemm_streamk/example_tilelang_gemm_streamk.py
- examples/gemv/example_gemv.py
- examples/grouped_gemm/example_grouped_gemm_bwd.py
- examples/grouped_gemm/example_grouped_gemm_fwd_ptr.py
- examples/kda/chunk_bwd_intra.py
- examples/kda/wy_fast_bwd.py
- examples/linear_attention/example_linear_attn_bwd.py
- examples/linear_attention/example_linear_attn_fwd.py
- examples/linear_attention/example_retention_fwd.py
- examples/minference/example_vertical_slash_sparse_attn.py
- examples/norm/rms_norm.py
- examples/norm/test_rms_norm.py
- examples/warp_specialize/example_warp_specialize_gemm_copy_1_gemm_0.py
- examples/warp_specialize/example_warp_specialize_gemm_copy_gemm_0_1.py
- maint/scripts/regression_all.py
- src/layout/gemm_layouts.cc
- src/layout/layout.cc
- src/layout/layout.h
- src/op/atomic_add.cc
- src/op/atomic_reduce.cc
- src/op/builtin.cc
- src/op/builtin.h
- src/op/copy.cc
- src/op/copy.h
- src/op/fill.cc
- src/op/finalize_reducer.cc
- src/op/gemm.cc
- src/op/gemm.h
- src/op/gemm_py.cc
- src/op/gemm_py.h
- src/op/gemm_sp.cc
- src/op/gemm_sp.h
- src/op/gemm_sp_py.cc
- src/op/gemm_sp_py.h
- src/op/operator.h
- src/op/reduce.cc
- src/op/reduce.h
- src/op/transpose.cc
- src/op/utils.cc
- src/op/utils.h
- src/target/codegen_cuda.cc
- src/target/codegen_cutedsl.cc
- src/target/codegen_hip.cc
- src/transform/common/pipeline_utils.h
- src/transform/common/tma_copy_utils.h
- src/transform/inject_pipeline.cc
- src/transform/instruction_annotation.cc
- src/transform/layout_inference.cc
- src/transform/legalize_safe_memory_access.cc
- src/transform/lower_ptx_async_copy.cc
- src/transform/lower_tile_op.cc
- src/transform/multi_version_buffer_rewriter.cc
- src/transform/multi_version_buffer_rewriter.h
- src/transform/optimize_cp_async_sync.cc
- src/transform/pipeline_planning.cc
- src/transform/producer_consumer_ws.cc
- src/transform/ptx_async_copy_injector.h
- src/transform/reuse_local_descriptor_allocations.cc
- testing/conftest.py
- testing/python/analysis/test_tilelang_nested_loop_checker.py
- testing/python/cache/test_tilelang_kernel_cache.py
- testing/python/carver/test_tilelang_carver_recommend_hints.py
- testing/python/components/test_tilelang_pass_config_disable_tma_lower.py
- testing/python/components/test_tilelang_pass_config_disable_warp_specialized.py
- testing/python/cuda/test_cuda_f32x2_intrinsics.py
- testing/python/issue/test_tilelang_issue_1001.py
- testing/python/issue/test_tilelang_issue_1008.py
- testing/python/issue/test_tilelang_issue_1106.py
- testing/python/issue/test_tilelang_issue_1210.py
- testing/python/issue/test_tilelang_issue_1263.py
- testing/python/issue/test_tilelang_issue_1744.py
- testing/python/issue/test_tilelang_issue_tma_no_ws.py
- testing/python/issue/test_tilelang_issue_ws_simt_copy_full_producer_extent.py
- testing/python/kernel/test_tilelang_kernel_bf16_gemm_tcgen5_ts.py
- testing/python/kernel/test_tilelang_kernel_gemm_batched.py
- testing/python/kernel/test_tilelang_kernel_gemm_with_stride.py
- testing/python/kernel/test_tilelang_kernel_int8_gemm_tcgen5.py
- testing/python/language/test_tilelang_language_all_of.py
- testing/python/language/test_tilelang_language_annotate_safe_value.py
- testing/python/language/test_tilelang_language_any_of.py
- testing/python/language/test_tilelang_language_chain_equal.py
- testing/python/language/test_tilelang_language_clear.py
- testing/python/language/test_tilelang_language_composable_index.py
- testing/python/language/test_tilelang_language_copy.py
artifact_dir: evidence/pull-bundles/tilelang/gh-2002
head_ref_oid: 144a9bd3a6bedd11ba85e307bcafdf916c25064b
---
# [Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS

- Repository: `tile-ai/tilelang`
- PR: [2002](https://github.com/tile-ai/tilelang/pull/2002)
- Decision: `Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.`
- Reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.

## Changed Paths

- `.gitignore`
- `examples/bitnet-1.58b/tokenization_bitnet.py`
- `examples/deepseek_mhc/example_mhc_bwd.py`
- `examples/deepseek_mhc/example_mhc_post.py`
- `examples/deepseek_mhc/example_mhc_pre.py`
- `examples/deepseek_mla/example_mla_decode.py`
- `examples/deepseek_nsa/benchmark/benchmark_nsa_fwd.py`
- `examples/deepseek_nsa/example_tilelang_nsa_bwd.py`
- `examples/deepseek_nsa/example_tilelang_nsa_decode.py`
- `examples/deepseek_nsa/example_tilelang_nsa_fwd.py`
- `examples/deepseek_nsa/example_tilelang_nsa_fwd_varlen.py`
- `examples/deepseek_v32/inference/convert.py`
- `examples/deepseek_v32/inference/kernel.py`
- `examples/deepseek_v32/sparse_mla_bwd.py`
- `examples/deepseek_v32/sparse_mla_fwd.py`
- `examples/dequantize_gemm/example_dequant_gemm_bf16_fp4_hopper.py`
- `examples/dsa_sparse_finetune/indexer_bwd.py`
- `examples/dsa_sparse_finetune/indexer_topk_reducesum.py`
- `examples/dsa_sparse_finetune/sparse_mla_bwd.py`
- `examples/dsa_sparse_finetune/sparse_mla_fwd.py`
- `examples/dsa_sparse_finetune/sparse_mla_topk_reducesum.py`
- `examples/flash_attention_sm100/gqa_bwd_bshd.py`
- `examples/flash_attention_sm100/gqa_fwd_bshd.py`
- `examples/flash_attention_sm100/mha_bwd_bshd.py`
- `examples/flash_attention_sm100/mha_fwd_bshd.py`
- `examples/flash_decoding/example_gqa_decode.py`
- `examples/fusedmoe/example_fusedmoe_tilelang.py`
- `examples/gdn/example_chunk_o_bwd.py`
- `examples/gdn/example_cumsum.py`
- `examples/gdn/example_wy_fast_bwd_split.py`
- `examples/gdn/test_example_gdn_compilation.py`
- `examples/gemm_fp8/example_tilelang_gemm_fp8_intrinsic.py`
- `examples/gemm_fp8/example_tilelang_gemm_fp8_sm100.py`
- `examples/gemm_sm100/README.md`
- `examples/gemm_sm100/gemm_mma.py`
- `examples/gemm_sm100/gemm_tcgen5mma.py`
- `examples/gemm_streamk/example_tilelang_gemm_streamk.py`
- `examples/gemv/example_gemv.py`
- `examples/grouped_gemm/example_grouped_gemm_bwd.py`
- `examples/grouped_gemm/example_grouped_gemm_fwd_ptr.py`
- `examples/kda/chunk_bwd_intra.py`
- `examples/kda/wy_fast_bwd.py`
- `examples/linear_attention/example_linear_attn_bwd.py`
- `examples/linear_attention/example_linear_attn_fwd.py`
- `examples/linear_attention/example_retention_fwd.py`
- `examples/minference/example_vertical_slash_sparse_attn.py`
- `examples/norm/rms_norm.py`
- `examples/norm/test_rms_norm.py`
- `examples/warp_specialize/example_warp_specialize_gemm_copy_1_gemm_0.py`
- `examples/warp_specialize/example_warp_specialize_gemm_copy_gemm_0_1.py`
- `maint/scripts/regression_all.py`
- `src/layout/gemm_layouts.cc`
- `src/layout/layout.cc`
- `src/layout/layout.h`
- `src/op/atomic_add.cc`
- `src/op/atomic_reduce.cc`
- `src/op/builtin.cc`
- `src/op/builtin.h`
- `src/op/copy.cc`
- `src/op/copy.h`
- `src/op/fill.cc`
- `src/op/finalize_reducer.cc`
- `src/op/gemm.cc`
- `src/op/gemm.h`
- `src/op/gemm_py.cc`
- `src/op/gemm_py.h`
- `src/op/gemm_sp.cc`
- `src/op/gemm_sp.h`
- `src/op/gemm_sp_py.cc`
- `src/op/gemm_sp_py.h`
- `src/op/operator.h`
- `src/op/reduce.cc`
- `src/op/reduce.h`
- `src/op/transpose.cc`
- `src/op/utils.cc`
- `src/op/utils.h`
- `src/target/codegen_cuda.cc`
- `src/target/codegen_cutedsl.cc`
- `src/target/codegen_hip.cc`
- `src/transform/common/pipeline_utils.h`

## Evidence Bundle

`evidence/pull-bundles/tilelang/gh-2002`
