Yongqi Zhuo fb1b11f1b8 [Feature] Warp specialization schedules and materialization (#2892) ...

This PR adds a new set of annotations for describing arbitrary warp specialization scheme, with the abstraction of pipelines. A pipeline supports 4 kinds of operations: producer_acquire, producer_commit, consumer_wait, and consumer_release. Behind the scenes, mbarriers are created to implement the synchronization between warps.

This PR introduces a new pass, MaterializeWSSchedule, that transforms an annotated kernel to a warp-specialized kernel.

The set of infrastructure enables further development of fully automatic warp specialization in TileLang. I am working on this.

2026-08-11 15:39:40 +08:00

..

test_nullable_buffer_params.py

[BugFix] ArgBinder: relax shared-shape binding for unused nullable buffers (#1870)

2026-02-24 11:43:55 +08:00

test_readonly_param_const_codegen.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_cluster_planning.py

[Backend] Add target detector/normalizer registeration and merge ExecutionBackend into backend module (#2323)

2026-06-09 00:35:07 +08:00

test_tilelang_transform_config_index_bitwidth.py

[Enhancement] Update examples and tests for improved type handling functionality (#1448)

2025-12-17 14:41:25 +08:00

test_tilelang_transform_decouple_type_cast.py

[BugFix] Skip DecoupleTypeCast on Evaluate roots to keep cp.async operands in their address spaces (#2869)

2026-08-10 13:59:29 +08:00

test_tilelang_transform_disable_memory_reuse.py

[CUDA][IR] Preserve dynamic shared memory aliases (#2240)

2026-05-21 18:12:58 +08:00

test_tilelang_transform_flatten_buffer.py

[TIR][IR] Update to use tirx (#2216)

2026-05-20 15:38:13 +08:00

test_tilelang_transform_fuse_mbarrier_arrive_expect_tx.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_hoist_broadcast_values.py

[Backend] Add target detector/normalizer registeration and merge ExecutionBackend into backend module (#2323)

2026-06-09 00:35:07 +08:00

test_tilelang_transform_if_stmt_binding.py

[BugFix] Preserve if condition evaluation during fan-out (#2764)

2026-07-25 14:30:13 +08:00

test_tilelang_transform_im2col_fallback.py

[Example] Migrate eligible examples to eager style (#2010)

2026-05-21 20:27:06 +08:00

test_tilelang_transform_inject_fence_proxy.py

[CUDA] Reduce template include overhead (#2474)

2026-07-01 16:53:07 +08:00

test_tilelang_transform_inject_set_max_nreg.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_Inject_software_pipeline.py

[BugFix][Transform] Never classify side-effecting binds as replayable (atomics re-executed at every use site since v0.1.11) (#2651)

2026-07-22 17:43:01 +08:00

test_tilelang_transform_inject_tcgen05_fence.py

[CUDA] Pack logical TMEM buffers into shared tcgen05.alloc arenas (#2831)

2026-08-03 13:53:12 +08:00

test_tilelang_transform_layout_inference.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_legalize_negative_index.py

[BugFix] Support runtime-dependent vector negative indices (#2654)

2026-07-16 18:30:46 +08:00

test_tilelang_transform_legalize_safe_memory_access.py

[BugFix] Support runtime-dependent vector negative indices (#2654)

2026-07-16 18:30:46 +08:00

test_tilelang_transform_legalize_vectorized_loop.py

[Enhancement] Update examples and tests for improved type handling functionality (#1448)

2025-12-17 14:41:25 +08:00

test_tilelang_transform_let_inline.py

[TIR][IR] Update to use tirx (#2216)

2026-05-20 15:38:13 +08:00

test_tilelang_transform_lexical_alloc_scope.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_loop_unswitching.py

[BugFix] Preserve loop steps during unswitching (#2741)

2026-07-23 14:41:19 +08:00

test_tilelang_transform_lower_device_kernel_launch.py

[BugFix] Use callee global symbols for cross-target calls (#2740)

2026-07-23 14:25:27 +08:00

test_tilelang_transform_lower_hopper_intrin.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_lower_ldgstg.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_lower_shared_barrier.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_lower_shared_tmem.py

[CUDA] Pack logical TMEM buffers into shared tcgen05.alloc arenas (#2831)

2026-08-03 13:53:12 +08:00

test_tilelang_transform_lower_tile_op.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_make_packed_api.py

[TIR][Runtime] Enforce host-evaluable assumptions at runtime (#2655)

2026-07-14 12:24:26 +08:00

test_tilelang_transform_materialize_ws_schedule.py

[Feature] Warp specialization schedules and materialization (#2892)

2026-08-11 15:39:40 +08:00

test_tilelang_transform_pipeline_barrier_ownership.py

[Example] Migrate eligible examples to eager style (#2010)

2026-05-21 20:27:06 +08:00

test_tilelang_transform_pipeline_planning.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_plan_update_buffer_allocation_location.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_preserve_loop_step.py

[BugFix] Correctly preserve loop step when unrolling loops (#2835)

2026-08-02 02:51:33 +08:00

test_tilelang_transform_producer_consumer_ws_persistent.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_producer_consumer_ws.py

[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite (#2863)

2026-08-04 18:22:40 +08:00

test_tilelang_transform_simplify.py

[Transform] Place auto WS producers in first warp group (#2315)

2026-06-04 14:47:52 +08:00

test_tilelang_transform_smem_swizzle_alignment.py

[CUDA] Align swizzled TMA shared buffers (#2391)

2026-06-15 01:42:06 +08:00

test_tilelang_transform_split_host_device.py

[TIR][Runtime] Enforce host-evaluable assumptions at runtime (#2655)

2026-07-14 12:24:26 +08:00

test_tilelang_transform_thread_sync.py

[BugFix] Fix two-instance modeling in ThreadSync cross-thread race checks (#2805)

2026-07-30 19:15:08 +08:00

test_tilelang_transform_tmem_arena_packing.py

[CUDA] Pack logical TMEM buffers into shared tcgen05.alloc arenas (#2831)

2026-08-03 13:53:12 +08:00

test_tilelang_transform_tmem_physical_shape.py

[CUDA] Pack logical TMEM buffers into shared tcgen05.alloc arenas (#2831)

2026-08-03 13:53:12 +08:00

test_tilelang_transform_unroll_loop.py

[Refactor] Recycle T.unroll(explicit=True) for early explicit unrolling (#2859)

2026-08-04 13:52:54 +08:00

test_tilelang_transform_verify_parallel_loop.py

[Enhancement] Disable data race check by default, opt-in via env var (#2851)

2026-08-03 16:55:27 +08:00