Chennes 0f13a2c548 [CUDA] Add 16-bit hpow and hfmod bridges (#3163) ...

CUDA lowering emits hpow and hfmod for 16-bit pow and fmod, but no CUDA
toolkit declares either name, so fp16/bf16 fmod and variable-exponent
fp16 pow fail to compile. Integer-exponent pow is unaffected: TileLang
expands it into repeated multiplication before lowering.

Add half_t and bfloat16_t overloads that evaluate with powf or fmodf in
float32 and convert back, matching the existing bridge pattern. Remove
math.h's hpow-to-powf macro: it shadows the new overload at the call
site, and compiles only where the result is assigned directly.

Add codegen and numerical coverage for fp16 pow and both fmod dtypes.
The bfloat16 hpow bridge completes the backend layer, while T.pow on
bfloat16 remains blocked by an upstream TVM dtype check.

2026-09-03 23:14:49 +08:00

..

test_tilelang_cast_rounding.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_intimm.py

[Lint] Phaseout Yapf format and embrace ruff format (#1417)

2025-12-12 17:22:18 +08:00

test_tilelang_language_access_ptr_codegen.py

[CUDA] Support int4 T.gemm (#2063)

2026-04-20 16:35:43 +08:00

test_tilelang_language_access_ptr.py

[BugFix] Preserve buffer element offsets in access pointers (#2727)

2026-07-22 15:42:39 +08:00

test_tilelang_language_alias.py

[Bugfix] Open Rocm ci test and fix some bugs. (#1443)

2026-01-06 02:18:29 +08:00

test_tilelang_language_all_of.py

[Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS (#2002)

2026-04-07 23:54:55 +08:00

test_tilelang_language_alloc_var_cmp.py

[Bugfix] Fix alloc_var re-bind warning when assigned with comparison ops (#1974)

2026-03-27 11:17:21 +08:00

test_tilelang_language_alloc.py

[Lang] Reject non-positive arrive_count in alloc_barrier and alloc_cluster_barrier (#3112)

2026-08-30 14:53:11 +08:00

test_tilelang_language_annot.py

[Bugfix] Open Rocm ci test and fix some bugs. (#1443)

2026-01-06 02:18:29 +08:00

test_tilelang_language_annotate_safe_value.py

[Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS (#2002)

2026-04-07 23:54:55 +08:00

test_tilelang_language_any_of.py

[Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS (#2002)

2026-04-07 23:54:55 +08:00

test_tilelang_language_assume.py

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

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

test_tilelang_language_async_copy_gemm_sm80.py

[Refactor] Improve cp.async lowering and add async_copy op (#1887)

2026-03-06 17:35:16 +08:00

test_tilelang_language_atom_mma.py

[Refactor] Remove intrinsic compatibility facade (#2812)

2026-07-30 14:26:21 +08:00

test_tilelang_language_atomic.py

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

2026-08-26 15:32:05 +08:00

test_tilelang_language_augassign.py

[Bugfix] Fix ast builder error for value -= 1 (#1825)

2026-02-09 17:30:22 +08:00

test_tilelang_language_ceildiv.py

[Feature] Expose multiple CUDA intrinsics (#2473)

2026-07-02 17:29:54 +08:00

test_tilelang_language_chain_equal.py

[Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS (#2002)

2026-04-07 23:54:55 +08:00

test_tilelang_language_clamp.py

[CUDA] Fix FP8 min/max codegen (#3047)

2026-08-18 17:50:27 +08:00

test_tilelang_language_clear.py

[BugFix] Fix make_int sign-extending negative int8 lanes (#2438)

2026-06-23 14:40:11 +08:00

test_tilelang_language_cluster.py

[Transform] Materialize kernel launch per backend (#2387)

2026-06-12 17:14:02 +08:00

test_tilelang_language_composable_index.py

[Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS (#2002)

2026-04-07 23:54:55 +08:00

test_tilelang_language_cooperative.py

[Fix] Add register to read A ptr in test_tilelang_language_cooperative.py (#1593)

2026-01-04 15:45:38 +08:00

test_tilelang_language_copy.py

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

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

test_tilelang_language_dialect.py

[TIR][Python] Split backend-specific op proxies (#2870)

2026-08-05 12:57:23 +08:00

test_tilelang_language_dtype_as_torch.py

[BugFix][CUDA] Support int4x2 and uint4x2 codegen (#3036)

2026-08-25 14:48:53 +08:00

test_tilelang_language_eager_ast_signature.py

Fix eager AST handling for *args and **kwargs (#2330)

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

test_tilelang_language_eager_jit.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_language_extract_lane.py

[TIR][Language] Add typed vector lane extraction API (#2789)

2026-07-28 17:18:25 +08:00

test_tilelang_language_fastmath.py

[CUDA] Add 16-bit overloads for CUTLASS fast-math functions (#3097)

2026-08-27 17:17:24 +08:00

test_tilelang_language_float4_e2m1_unpacked_dtype.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_language_frontend_v2.py

[Testing] Run portable regression tests on auto targets (#2914)

2026-08-07 15:07:26 +08:00

test_tilelang_language_func_attrs.py

[JIT][FFI] Allocate kernel outputs through the packed API (#2937)

2026-08-11 10:47:08 +08:00

test_tilelang_language_get_warp_info.py

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

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

test_tilelang_language_if_range.py

[Refactor] Remove GEMM v1 and promote gemm_py to be the canonical gemm op (#2033)

2026-04-13 12:42:44 +08:00

test_tilelang_language_infinity.py

[BugFix] Support T.infinity for float8_e5m2 (#2671)

2026-07-15 17:37:20 +08:00

test_tilelang_language_int64.py

[Feature] Add full Windows support and fix related cross-platform issues (#2093)

2026-05-06 23:08:24 +08:00

test_tilelang_language_intrinsics_codegen.py

[CUDA] Add 16-bit hpow and hfmod bridges (#3163)

2026-09-03 23:14:49 +08:00

test_tilelang_language_isfinite_codegen.py

[Feature] Add CUDA intrinsic for isfinite operation (#1925)

2026-03-12 17:21:01 +08:00

test_tilelang_language_kernel_threads.py

[BugFix] Reject non-positive thread extents in T.Kernel (#2653)

2026-07-13 17:29:01 +08:00

test_tilelang_language_ldg.py

[Refactor] Unify the usage of cast-related operators (#1757)

2026-02-02 18:47:18 +08:00

test_tilelang_language_ldst.py

[Feature] Expose multiple CUDA intrinsics (#2473)

2026-07-02 17:29:54 +08:00

test_tilelang_language_let_layout.py

[Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS (#2002)

2026-04-07 23:54:55 +08:00

test_tilelang_language_let.py

[Fix] Bind statement missing role marker (#2362)

2026-06-12 17:35:42 +08:00

test_tilelang_language_mask_op.py

[Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS (#2002)

2026-04-07 23:54:55 +08:00

test_tilelang_language_min_blocks_per_sm.py

[Refactor] Remove GEMM v1 and promote gemm_py to be the canonical gemm op (#2033)

2026-04-13 12:42:44 +08:00

test_tilelang_language_negative_index.py

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

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

test_tilelang_language_nvf4_mma_block_scale.py

[CUDA] Add SM120 NVF4 block-scale MMA support (#2364)

2026-07-30 18:10:10 +08:00

test_tilelang_language_parallel.py

[BugFix] Always vectorize T.Parallel loops (#3121)

2026-08-31 23:33:42 +08:00

test_tilelang_language_pdl.py

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

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

test_tilelang_language_pipeline.py

[Refactor][Pipeline] Run pipeline rewriting before layout inference and stabilize tiled WS (#2002)

2026-04-07 23:54:55 +08:00

test_tilelang_language_ptr.py

[BugFix] Skip descriptor TMA for device-bound copy bases (#2803)

2026-07-29 20:35:07 +08:00

test_tilelang_language_rand.py

[CUDA][Codegen] Keep RNG state in kernel scope (#2540)

2026-07-07 14:41:03 +08:00

test_tilelang_language_reduce_maxmin_nan.py

[BugFix] Honor nan_propagate in reduce max/min/absmax clear=False write-back (#2788)

2026-07-28 16:04:07 +08:00

test_tilelang_language_reduce.py

[Lang][Reducer] Allow conditional reducer finalization (#3044)

2026-08-17 16:59:29 +08:00

test_tilelang_language_reducer_v2.py

[Refactor] Give reducers a first-class PartialFragment layout solved by layout inference (#3093)

2026-08-27 14:20:55 +08:00

test_tilelang_language_reshape.py

[Refactor] Move backend-specific GEMM implementations and transforms into backend directories (#2165)

2026-05-08 17:11:34 +08:00

test_tilelang_language_scan.py

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

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

test_tilelang_language_select.py

[BugFix] Scalarize Select in automatic vectorization (#3060)

2026-08-24 12:43:44 +08:00

test_tilelang_language_source_kernel.py

[Feature] Introduce T.CUDASourceCodeKernel (#1970)

2026-04-17 16:48:28 +08:00

test_tilelang_language_span.py

[TIR] Inject source spans into tirx IR and surface source locations in compiler errors (#2751)

2026-07-28 13:46:14 +08:00

test_tilelang_language_subtype.py

[Bugfix][Subtype] Fix scalar fp4 store/load codegen for non-packed buffers (#2037)

2026-04-14 15:54:28 +08:00

test_tilelang_language_sync_threads.py

Allow variable barrier id in T.sync_threads() (#2197)

2026-05-18 14:50:43 +08:00

test_tilelang_language_tcgen05_cutedsl.py

[Layout] Support half-subpartition (M=64) TMEM tiles in tcgen05.ld/st (#2880)

2026-08-11 15:35:15 +08:00

test_tilelang_language_tcgen05_gemm.py

[Bugfix] Fix wait parity for explicit mbarriers in pipelined loops (#3087)

2026-08-26 15:18:10 +08:00

test_tilelang_language_tcgen05_sliced_operands.py

[CUDA] Support arbitrary TMEM layouts (#2785)

2026-07-29 13:07:28 +08:00

test_tilelang_language_ternary.py

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

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

test_tilelang_language_tile_schedule.py

[Language][Scheduler] Expose scalar tile scheduler state (#2553)

2026-07-08 14:34:55 +08:00

test_tilelang_language_tma_1d.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_language_tma_copy.py

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

2026-08-29 15:43:29 +08:00

test_tilelang_language_tma_gather_scatter.py

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

2026-08-29 15:43:29 +08:00

test_tilelang_language_tma_store.py

[Feature] Add read option to TMA store wait (#2300)

2026-05-31 03:51:36 +08:00

test_tilelang_language_tmem_copy.py

[BugFix][CUDA] Advance tcgen05 ld/st segment pointers in b32 columns (#2952)

2026-08-13 13:35:42 +08:00

test_tilelang_language_transpose.py

[Feature] Add full Windows support and fix related cross-platform issues (#2093)

2026-05-06 23:08:24 +08:00

test_tilelang_language_unroll.py

[Enhancement] Improve unroll loop functionality for dynamic extent and corresponding test case (#1654)

2026-01-12 22:33:17 +08:00

test_tilelang_language_var_init.py

[Testing] Run portable regression tests on auto targets (#2914)

2026-08-07 15:07:26 +08:00

test_tilelang_language_vectorize_matrix.py

[Refactor][CUDA] Replace per-lane-count vector constructors with variadic packers (#3158)

2026-09-03 23:01:35 +08:00

test_tilelang_language_vectorize.py

[BugFix][CUDA] Support non-constant int4/uint4 broadcast (#3114)

2026-09-02 16:45:13 +08:00

test_tilelang_language_vectorized_cast.py

[Enhancement] Add vectorized fp8x2 <-> fp16/bf16 cast codegen (#2475)

2026-06-29 16:00:57 +08:00

test_tilelang_language_view.py

[Layout][CUDA] Support reinterpreting (dtype-changing) T.view aliases (#2953)

2026-08-13 13:01:51 +08:00

test_tilelang_language_warp_reduce.py

[BugFix] Fix warp_reduce truncating int64/uint64 to 32 bits on sm_80+ (#2782)

2026-07-29 16:39:26 +08:00

test_tilelang_language_warp_sync.py

[ROCm] Preserve FP8 bits in warp shuffles (#3104)

2026-08-29 13:27:01 +08:00

test_tilelang_language_warp_vote.py

[BugFix] Reject floating-point predicates in vote intrinsics (#2800)

2026-07-29 13:33:48 +08:00

test_tilelang_language_wgmma_gemm.py

[CUDA] Adopt multi-staged buffers in examples (#2836)

2026-08-04 11:26:00 +08:00

test_tilelang_memory_leak.py

[Cherry-Pick][BugFix] Fix get_parent_locals frame self-reference leak (#2934)

2026-08-10 16:08:38 +08:00