---
id: pr-tilelang-1840
repo: tile-ai/tilelang
pr: 1840
title: '[BugFix] Fix Hopper TMA lowering without warp specialization'
author: Henry-Jessie
date: '2026-02-25'
url: https://github.com/tile-ai/tilelang/pull/1840
source_category: upstream-code
architectures:
- sm90
tags:
- TMA
- WGMMA
- arch_pipeline
- copy lowering
- hopper
techniques:
- async-copy
- pipeline-stages
- warp-specialization
hardware_features:
- hopper
- tma
- wgmma
kernel_types:
- runtime
languages:
- cute-dsl
captured_at: '2026-05-16'
status: merged
merge_sha: 911385fdd12238209edb410be4ba2c722a92d42a
inclusion_reason: 'Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.'
changed_paths:
- src/transform/inject_tma_barrier.cc
- testing/python/issue/test_tilelang_issue_tma_no_ws.py
- tilelang/engine/phase.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1840
head_ref_oid: d24ef1c2aaecea340c3ee3a1843a4b1ba751333c
---
# [BugFix] Fix Hopper TMA lowering without warp specialization

- Repository: `tile-ai/tilelang`
- PR: [1840](https://github.com/tile-ai/tilelang/pull/1840)
- 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

- `src/transform/inject_tma_barrier.cc`
- `testing/python/issue/test_tilelang_issue_tma_no_ws.py`
- `tilelang/engine/phase.py`

## Evidence Bundle

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