mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-25 07:14:16 +08:00
27f1f81c61
* [Feature] Add Blackwell Cluster Launch Control (CLC) primitives
Expose cluster launch control query operations (try_cancel,
try_cancel_multicast, is_canceled, get_first_ctaid_{x,y,z}) in both
the C++ codegen path and the Python tilelang.language API, plus a new
mbarrier_arrive_expect_tx builtin used with CLC multicast completion.
Enables writing persistent SM100 kernels that use the hardware scheduler
to fetch the next tile coordinate.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* [Example] Add SM100 persistent 2-CTA GEMM example using CLC scheduling
Adds examples/gemm_sm100/gemm_tcgen5mma_ws_clc.py, a warp-specialized
persistent GEMM for Blackwell (sm_100) that uses tcgen05 2-CTA MMA
with cluster launch control tile scheduling. The kernel is split into
producer/MMA/scheduler/epilogue warp groups, with the scheduler warp
running on both CTAs and feeding next-tile coordinates via the CLC
hardware multicast into schedule_arrived directly (no intermediate
barrier). Storage sync is disabled and all syncs are placed manually.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* lint
* typo
Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>
* remove unnecessary test
* Add CLC query intrinsics and async fence
* Refactor CLC response handling by introducing CLCResponseDecode struct. Update clc_get_first_ctaid_x/y/z and clc_is_canceled functions to utilize the new decode response method for improved clarity and maintainability.
* Enhance CLC scheduling in persistent GEMM by introducing shared variables for schedule validity and tile ID. Update cancellation checks to utilize the new shared variables, improving clarity and maintainability of the scheduling logic.
* test fix
---------
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>
Co-authored-by: Zhiwen Mo <zm125@ic.ac.uk>
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>