---
id: pr-tilelang-1909
repo: tile-ai/tilelang
pr: 1909
title: '[Feature] Add Producer-Consumer Warp Specialization and T.tma_copy() API'
author: LeiWang1999
date: '2026-03-18'
url: https://github.com/tile-ai/tilelang/pull/1909
source_category: upstream-code
architectures:
- sm100
- sm90
tags:
- MLA
- TMA
- WGMMA
- arch_pipeline
- attention_kv
- benchmark_test
- copy lowering
- gemm_quant
- hopper
- memory_primitives
- mxfp4
techniques:
- async-copy
- fine-grained-quantization
- pipeline-stages
- warp-specialization
hardware_features:
- fp4
- hopper
- mxfp4
- tma
- wgmma
kernel_types:
- attention
- decode
- gemm
- memory
- mla
- runtime
languages:
- cute-dsl
captured_at: '2026-05-16'
status: merged
merge_sha: ded6a992218a5ec0dbb490128cb21aae7794cdb8
inclusion_reason: 'Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.'
changed_paths:
- examples/attention_sink/benchmark_gqa_sink_fwd.py
- examples/attention_sink/benchmark_mha_sink_fwd.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_fwd_bhsd_wgmma_pipelined.py
- examples/attention_sink/regression_attention_sink.py
- examples/attention_sink/test_example_attention_sink.py
- examples/blocksparse_attention/example_tilelang_sparse_gqa_decode_varlen_indice.py
- examples/blocksparse_gemm/example_blocksparse_gemm.py
- examples/deepseek_v32/sparse_mla_fwd_pipelined.py
- examples/dequantize_gemm/example_dequant_gemm_w4a8.py
- examples/dequantize_gemm/example_dequant_groupedgemm_bf16_mxfp4_hopper.py
- examples/flash_attention/example_gqa_bwd_wgmma_pipelined.py
- examples/flash_attention/example_gqa_fwd_bshd_wgmma_pipelined.py
- examples/flash_attention/example_mha_bwd_bshd_wgmma_pipelined.py
- examples/flash_attention/example_mha_fwd_bhsd_wgmma_pipelined.py
- examples/flash_attention/test_example_flash_attention.py
- examples/gdn/test_example_gdn_compilation.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1909
head_ref_oid: 61449a9c0e444f093982b30b79024e42cdbbee93
---
# [Feature] Add Producer-Consumer Warp Specialization and T.tma_copy() API

- Repository: `tile-ai/tilelang`
- PR: [1909](https://github.com/tile-ai/tilelang/pull/1909)
- Decision: `Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.`
- Reason: Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.

## Changed Paths

- `examples/attention_sink/benchmark_gqa_sink_fwd.py`
- `examples/attention_sink/benchmark_mha_sink_fwd.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_fwd_bhsd_wgmma_pipelined.py`
- `examples/attention_sink/regression_attention_sink.py`
- `examples/attention_sink/test_example_attention_sink.py`
- `examples/blocksparse_attention/example_tilelang_sparse_gqa_decode_varlen_indice.py`
- `examples/blocksparse_gemm/example_blocksparse_gemm.py`
- `examples/deepseek_v32/sparse_mla_fwd_pipelined.py`
- `examples/dequantize_gemm/example_dequant_gemm_w4a8.py`
- `examples/dequantize_gemm/example_dequant_groupedgemm_bf16_mxfp4_hopper.py`
- `examples/flash_attention/example_gqa_bwd_wgmma_pipelined.py`
- `examples/flash_attention/example_gqa_fwd_bshd_wgmma_pipelined.py`
- `examples/flash_attention/example_mha_bwd_bshd_wgmma_pipelined.py`
- `examples/flash_attention/example_mha_fwd_bhsd_wgmma_pipelined.py`
- `examples/flash_attention/test_example_flash_attention.py`
- `examples/gdn/test_example_gdn_compilation.py`

## Evidence Bundle

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