---
id: pr-tilelang-2003
repo: tile-ai/tilelang
pr: 2003
title: '[Transform] Add InjectTcgen05Fence pass'
author: LeiWang1999
date: '2026-04-16'
url: https://github.com/tile-ai/tilelang/pull/2003
source_category: upstream-code
architectures:
- sm100
tags:
- arch_pipeline
- blackwell
- compiler_runtime
- copy lowering
- cuda backend
- intrinsic
techniques:
- persistent-kernels
- pipeline-stages
hardware_features:
- blackwell
- tcgen05
- tmem
kernel_types:
- compiler
- gemm
- runtime
languages:
- cuda-cpp
- cute-dsl
captured_at: '2026-05-16'
status: merged
merge_sha: 891109ea879839bbfe2ad937db208f2d7c6f1ce0
inclusion_reason: 'Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.'
changed_paths:
- examples/gemm_sm100/README.md
- examples/gemm_sm100/gemm_tcgen5mma_ws.py
- examples/gemm_sm100/gemm_tcgen5mma_ws_persistent.py
- src/op/builtin.cc
- src/op/builtin.h
- src/op/copy.cc
- src/target/codegen_cuda.cc
- src/target/codegen_cutedsl.cc
- src/transform/inject_tcgen05_fence.cc
- testing/python/transform/test_tilelang_transform_inject_tcgen05_fence.py
- tilelang/engine/phase.py
- tilelang/transform/__init__.py
artifact_dir: evidence/pull-bundles/tilelang/gh-2003
head_ref_oid: dfbe9a19a086e198d69701977e1fd1e4e614aa1f
---
# [Transform] Add InjectTcgen05Fence pass

- Repository: `tile-ai/tilelang`
- PR: [2003](https://github.com/tile-ai/tilelang/pull/2003)
- 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/README.md`
- `examples/gemm_sm100/gemm_tcgen5mma_ws.py`
- `examples/gemm_sm100/gemm_tcgen5mma_ws_persistent.py`
- `src/op/builtin.cc`
- `src/op/builtin.h`
- `src/op/copy.cc`
- `src/target/codegen_cuda.cc`
- `src/target/codegen_cutedsl.cc`
- `src/transform/inject_tcgen05_fence.cc`
- `testing/python/transform/test_tilelang_transform_inject_tcgen05_fence.py`
- `tilelang/engine/phase.py`
- `tilelang/transform/__init__.py`

## Evidence Bundle

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