loopyt 39b086d781 feat(h20): route tile_kernels calls through a compat layer, so the model runs without TileLang ... The references were importable and numerically validated, but nothing pointed the model at them: the forward pass still called the compiled kernels, which raise ModuleNotFoundError on a host without TileLang. This adds the routing layer and swaps the call sites. `deepinfer/tk_compat/` exposes the *compiled* signatures -- out-buffers, in-place conventions, argument order -- because that is what the 25 call sites are written against, and adapts to the reference's shape internally. Selection is `DEEPINFER_KERNELS=auto|cuda|ref`; `auto` follows whether TileLang imports. The adapters are not pass-throughs. Each difference was found by running both backends: * `per_token_cast` takes `num_per_channels=n`, the reference takes `block_size=(1, n)`. * `apply_rotary` takes a concatenated cos|sin cache and rotates in place returning None; the reference takes cos and sin separately and returns a new tensor. The adapter splits, calls, and copies back, because the call site relies on the in-place contract. * `engram_gate` takes the pre-multiplied weight and `(eps, clamp_value)`; the reference needs the two weights separately and `(clamp_value, eps)`. Since a product cannot be un-multiplied, `fused_weight` returns the pair on the reference path and the gate unpacks it. * `mhc_pre_big_fuse_fwd` has no single reference: its five steps are composed from the reference's own components. * `mhc_reduce_partials_and_rmsnorm_fwd` likewise; rebuilt from the reference's formula. `moe_topk_gate(force_random=...)` raises rather than silently ignoring the flag. The reference has no equivalent, and dropping it would route tokens differently while looking healthy. Verification, all on the real call signatures: 18 cases compared between `cuda` and `ref` modes, `ALL_MATCH`, 12 bitwise identical including every quantised byte and every routing index; plus a smoke run that executes all 23 APIs with TileLang blocked, `ALL_CALLS_OK`. End to end still has not completed, and the reason is not this change: six GPUs on the host are held by other containers (82 GiB of 139 each), so the 8-way EP run OOMs in allocate_userbuffer and the surviving ranks wait at a barrier. Recorded in the doc along with the command to rerun it.		2026-10-22 18:10:14 +08:00
..
simulate_no_tilelang	feat(h20): make the pure-torch kernel references usable without TileLang	2026-10-22 16:08:46 +08:00
tile_kernels_patch	feat(h20): route tile_kernels calls through a compat layer, so the model runs without TileLang	2026-10-22 18:10:14 +08:00
ab_cases.py	test(h20): catch a factory that returns nothing before eight ranks do	2026-10-22 14:16:41 +08:00
build_vendor.sh	initial commit	2026-10-20 12:09:42 +08:00
build.sh	initial commit	2026-10-20 12:09:42 +08:00
compare_layers.py	fix(sm90): expand expert scales on the contraction axis only	2026-10-21 16:50:31 +08:00
copy_site_packages.py	feat(h20): make the FlagTree image setup reproducible instead of a shell history	2026-10-22 13:20:03 +08:00
enable_torch_kernel_refs.sh	feat(h20): make the pure-torch kernel references usable without TileLang	2026-10-22 16:08:46 +08:00
flagtree_triton_probe.py	docs(h20): record what a FlagTree/FlagOS image can and cannot do	2026-10-22 10:53:47 +08:00
format_csrc.sh	initial commit	2026-10-20 12:09:42 +08:00
gen_stub.sh	initial commit	2026-10-20 12:09:42 +08:00
h20_ask.py	style(sm90): satisfy ruff and add the missing type annotations	2026-10-21 17:09:08 +08:00
h20_attn_check.py	style(sm90): satisfy ruff and add the missing type annotations	2026-10-21 17:09:08 +08:00
h20_e2e_report.sh	docs(h20): point at the setup script, and let the report script carry its own defaults	2026-10-22 13:20:08 +08:00
h20_engram_check.py	style(sm90): satisfy ruff and add the missing type annotations	2026-10-21 17:09:08 +08:00
h20_expert_weight_check.py	style(sm90): satisfy ruff and add the missing type annotations	2026-10-21 17:09:08 +08:00
h20_index_decode_check.py	test(sm90): rule out the KV cache and the decode indexer	2026-10-21 20:02:19 +08:00
h20_kv_roundtrip.py	test(sm90): rule out the KV cache and the decode indexer	2026-10-21 20:02:19 +08:00
h20_ops_ab.py	test(h20): catch a factory that returns nothing before eight ranks do	2026-10-22 14:16:41 +08:00
h20_philox_check.py	feat(triton): reproduce cuRAND's Philox stream, so the sampler can be ported	2026-10-22 08:37:08 +08:00
h20_quant_check.py	test(sm90): decode FP4 by table lookup in the quant check	2026-10-21 17:15:16 +08:00
h20_real_requant_check.py	style(sm90): satisfy ruff and add the missing type annotations	2026-10-21 17:09:08 +08:00
h20_recipe_probe.py	debug: capture final logits, and keep the profile forward out of the dumps	2026-10-21 17:44:21 +08:00
h20_rng_dump.py	feat(triton): reproduce cuRAND's Philox stream, so the sampler can be ported	2026-10-22 08:37:08 +08:00
h20_smoke_check.py	style(sm90): satisfy ruff and add the missing type annotations	2026-10-21 17:09:08 +08:00
h20_state_check.py	test(h20): check what the model actually says, in a file rather than on the command line	2026-10-22 10:53:01 +08:00
h20_topk_check.py	style(sm90): satisfy ruff and add the missing type annotations	2026-10-21 17:09:08 +08:00
ref_layer_dump.py	debug: capture final logits, and keep the profile forward out of the dumps	2026-10-21 17:44:21 +08:00
ref_next_token.py	debug: per-step decode dumps, and prove the flip is decode-only	2026-10-21 19:07:49 +08:00
setup_flagtree_image.sh	feat(h20): make the FlagTree image setup reproducible instead of a shell history	2026-10-22 13:20:03 +08:00
verify_tile_kernels_patch.py	feat(h20): make the pure-torch kernel references usable without TileLang	2026-10-22 16:08:46 +08:00