---
id: pr-tilelang-1882
repo: tile-ai/tilelang
pr: 1882
title: '[Feature] 2-SM support for TMA, TMEM and TCGEN5MMA on Blackwell'
author: Rachmanino
date: '2026-03-24'
url: https://github.com/tile-ai/tilelang/pull/1882
source_category: upstream-code
architectures:
- sm100
tags:
- TMA
- arch_pipeline
- benchmark
- blackwell
- copy lowering
- tcgen
techniques:
- async-copy
- persistent-kernels
- pipeline-stages
- swizzling
- warp-specialization
hardware_features:
- blackwell
- tcgen05
- tma
- tmem
kernel_types:
- gemm
- runtime
languages:
- cuda-cpp
- cute-dsl
captured_at: '2026-05-16'
status: merged
merge_sha: dd0cd3ebc9cb568a0ec367bcf06f190a8762c69a
inclusion_reason: 'Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.'
changed_paths:
- examples/gemm_sm100/gemm_tcgen5mma_ws.py
- examples/gemm_sm100/gemm_tcgen5mma_ws_persistent.py
- maint/gemm_v2/correctness_evaluation_tcgen05.py
- maint/gemm_v2/correctness_evaluation_tcgen05_2cta.py
- src/op/builtin.h
- src/op/copy.cc
- src/op/gemm_py.cc
- src/op/gemm_py.h
- src/op/operator.h
- src/op/tcgen5_meta.h
- src/target/codegen_cuda.cc
- src/target/codegen_cuda.h
- src/target/codegen_cutedsl.cc
- src/tl_templates/cuda/cluster.h
- src/tl_templates/cuda/common.h
- src/tl_templates/cuda/copy_sm100.h
- src/transform/lower_shared_tmem.cc
- src/transform/lower_tile_op.cc
- tilelang/engine/phase.py
- src/tl_templates/cuda/instruction/tcgen05mma.h
- src/tl_templates/cuda/tcgen_05.h
- src/tl_templates/cuda/threadblock_swizzle.h
- src/transform/fuse_mbarrier_arrive_expect_tx.cc
- src/transform/lower_blackwell_2sm.cc
artifact_dir: evidence/pull-bundles/tilelang/gh-1882
head_ref_oid: 28f4224a0d42562801f92e9eeb5a9508110e31d3
---
# [Feature] 2-SM support for TMA, TMEM and TCGEN5MMA on Blackwell

- Repository: `tile-ai/tilelang`
- PR: [1882](https://github.com/tile-ai/tilelang/pull/1882)
- 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/gemm_sm100/gemm_tcgen5mma_ws.py`
- `examples/gemm_sm100/gemm_tcgen5mma_ws_persistent.py`
- `maint/gemm_v2/correctness_evaluation_tcgen05.py`
- `maint/gemm_v2/correctness_evaluation_tcgen05_2cta.py`
- `src/op/builtin.h`
- `src/op/copy.cc`
- `src/op/gemm_py.cc`
- `src/op/gemm_py.h`
- `src/op/operator.h`
- `src/op/tcgen5_meta.h`
- `src/target/codegen_cuda.cc`
- `src/target/codegen_cuda.h`
- `src/target/codegen_cutedsl.cc`
- `src/tl_templates/cuda/cluster.h`
- `src/tl_templates/cuda/common.h`
- `src/tl_templates/cuda/copy_sm100.h`
- `src/transform/lower_shared_tmem.cc`
- `src/transform/lower_tile_op.cc`
- `tilelang/engine/phase.py`
- `src/tl_templates/cuda/instruction/tcgen05mma.h`
- `src/tl_templates/cuda/tcgen_05.h`
- `src/tl_templates/cuda/threadblock_swizzle.h`
- `src/transform/fuse_mbarrier_arrive_expect_tx.cc`
- `src/transform/lower_blackwell_2sm.cc`

## Evidence Bundle

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