---
id: pr-tilelang-1874
repo: tile-ai/tilelang
pr: 1874
title: '[Feature] Support cluster launch, query, synchronization and barrier operations'
author: Rachmanino
date: '2026-02-28'
url: https://github.com/tile-ai/tilelang/pull/1874
source_category: upstream-code
architectures:
- sm100
- sm90
tags:
- WGMMA
- arch_pipeline
- blackwell
- compiler_runtime
- cuda backend
- scheduler_autotune
techniques:
- autotuning
- persistent-kernels
- pipeline-stages
- swizzling
- tile-scheduling
hardware_features:
- blackwell
- hopper
- wgmma
kernel_types:
- compiler
- gemm
- runtime
- scheduler
languages:
- cuda-cpp
- cute-dsl
captured_at: '2026-05-16'
status: merged
merge_sha: 2e854b068eedcba54593d298481247131ecd7829
inclusion_reason: 'Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.'
changed_paths:
- 3rdparty/tvm
- examples/gemm_sm100/gemm_tcgen5mma_ws_persistent.py
- src/layout/swizzle.cc
- src/layout/swizzle.h
- src/op/builtin.cc
- src/op/builtin.h
- src/target/codegen_cuda.cc
- src/target/codegen_cuda.h
- src/target/codegen_cutedsl.cc
- src/target/rt_mod_cuda.cc
- src/tl_templates/cuda/cluster.h
- src/tl_templates/cuda/tcgen_05.h
- src/transform/eliminate_storage_sync_for_mbarrier.cc
- src/transform/lower_device_kernel_launch.cc
- src/transform/lower_device_storage_access_info.cc
- src/transform/lower_opaque_block.cc
- src/transform/lower_shared_barrier.cc
- tilelang/engine/phase.py
- tilelang/jit/adapter/wrapper.py
- tilelang/language/__init__.py
- tilelang/language/allocate.py
- tilelang/language/builtin.py
- tilelang/language/cluster.py
- tilelang/language/kernel.py
- testing/python/language/test_tilelang_language_cluster.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1874
head_ref_oid: de0b7ff67ea22e37c5427898d12d0e13259723d6
---
# [Feature] Support cluster launch, query, synchronization and barrier operations

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

- `3rdparty/tvm`
- `examples/gemm_sm100/gemm_tcgen5mma_ws_persistent.py`
- `src/layout/swizzle.cc`
- `src/layout/swizzle.h`
- `src/op/builtin.cc`
- `src/op/builtin.h`
- `src/target/codegen_cuda.cc`
- `src/target/codegen_cuda.h`
- `src/target/codegen_cutedsl.cc`
- `src/target/rt_mod_cuda.cc`
- `src/tl_templates/cuda/cluster.h`
- `src/tl_templates/cuda/tcgen_05.h`
- `src/transform/eliminate_storage_sync_for_mbarrier.cc`
- `src/transform/lower_device_kernel_launch.cc`
- `src/transform/lower_device_storage_access_info.cc`
- `src/transform/lower_opaque_block.cc`
- `src/transform/lower_shared_barrier.cc`
- `tilelang/engine/phase.py`
- `tilelang/jit/adapter/wrapper.py`
- `tilelang/language/__init__.py`
- `tilelang/language/allocate.py`
- `tilelang/language/builtin.py`
- `tilelang/language/cluster.py`
- `tilelang/language/kernel.py`
- `testing/python/language/test_tilelang_language_cluster.py`

## Evidence Bundle

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