Yongqi ZhuoandClaude Fable 5 30c1904d3a [CUDA] Unify TMA copy lowering on CuTe algebra (#3106) ...

* [CUDA] Unify TMA copy lowering on CuTe algebra and support general swizzled layouts

Every descriptor-based TMA path (bulk copy, atomic add, gather4/scatter4,
im2col) now shares one CuTe-algebra box decomposition. Boxes follow the
driver contract (inner box bytes at most the swizzle span, per-dim caps),
and the recovered swizzle is widened on demand under an equivalence proof,
so general annotated layouts -- including phase-shifted swizzles -- lower
correctly.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

* Address review concerns

---------

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>

2026-08-29 15:43:29 +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_codegen_line_directives.py

[CodeGen][ROCm] Emit #line directives for HIP target (#3058)

2026-08-24 12:44:19 +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

[CUDA] Unify TMA copy lowering on CuTe algebra (#3106)

2026-08-29 15:43:29 +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] Fix wait parity for explicit mbarriers in pipelined loops (#3087)

2026-08-26 15:18:10 +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

[Layout][Reducer] Preserve vectorized reducer update plans (#3100)

2026-08-27 22:20:13 +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_loop_vectorize.py

[BugFix][Transform] Avoid int32 overflow in vector analysis (#3066)

2026-08-27 19:53:47 +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

[Layout] Reject non-round-tripping inferred inverses (#3090)

2026-08-27 18:04:12 +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

[Bugfix] Make TMA layouts region-aware to keep slices contiguous (#3089)

2026-08-26 15:32:05 +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_span_preservation.py

[TIR][Transform] Preserve source spans across lowering passes (#2966)

2026-08-17 15:51:43 +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

[TIR][Transform] Handle grid barriers and unbounded pointer ranges (#3050)

2026-08-19 14:30:50 +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_buffer_init.py

[Transform] Add VerifyBufferInit, a general buffer-initialization check (#2956)

2026-08-17 12:12:36 +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