Skip to content

[Feature] Support cluster launch, query, synchronization and barrier operations #1874

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
LeiWang1999 merged 13 commits into from
Feb 28, 2026
Merged

[Feature] Support cluster launch, query, synchronization and barrier operations #1874

Changes from 1 commit
Commits
Show all changes
13 commits
Select commit Hold shift + click to select a range
e22e876
support cluster launch for tvm_ffi and cython backend
Rachmanino Feb 24, 2026
2f0e642
lint
Rachmanino Feb 24, 2026
e5c1ec7
add test
Rachmanino Feb 25, 2026
7c3a425
lint
Rachmanino Feb 25, 2026
13e3d20
update tvm
Rachmanino Feb 26, 2026
acfaf3f
Add cluster synchronization operations and related functions
Rachmanino Feb 26, 2026
f5c4848
add test for intrinsics
Rachmanino Feb 26, 2026
c8612dc
lint
Rachmanino Feb 26, 2026
8aa3a2b
remove legacy code
Rachmanino Feb 26, 2026
edd8f50
Add cluster barrier support and related functionality
Rachmanino Feb 26, 2026
e94be2c
lint
Rachmanino Feb 26, 2026
1bcc3a3
fix typo
Rachmanino Feb 26, 2026
de0b7ff
add storage sync for cluster
Rachmanino Feb 26, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
lint
  • Loading branch information
@Rachmanino
Rachmanino committed Feb 26, 2026
commit 7c3a4258b0feb302d642a90a085b3b3d4926f242
5 changes: 2 additions & 3 deletions testing/python/language/test_tilelang_language_cluster_launch.py
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,5 @@
import tilelang
import tilelang.language as T
import torch
import tilelang.testing


Expand Down Expand Up @@ -30,7 +29,7 @@ def gemm(
def run_cython_cluster_launch():
kernel = matmul(1024, 1024, 1024, 128, 128, 32)
mod = tilelang.compile(kernel, execution_backend="cython")
assert 'clusterDim = {2, 1, 1}' in mod.get_host_source()
assert "clusterDim = {2, 1, 1}" in mod.get_host_source()


def run_tvm_ffi_cluster_launch():
Expand All @@ -48,7 +47,7 @@ def run_tvm_ffi_cluster_launch():
(((TVMFFIAny*)stack_ffi_any)[5].v_int64) = ((int64_t)1);
"""
assert check_str in mod.get_host_source()


@tilelang.testing.requires_cuda
@tilelang.testing.requires_cuda_compute_version_ge(9, 0)
Expand Down