Section 6.0 now states the status as a table, because the previous wording invited a wrong
reading. Importing the torch references -- and even deepinfer itself -- succeeds without
TileLang, and five operators are numerically validated. But the model forward still calls the
compiled quantisation and MoE entry points, which raise ModuleNotFoundError, because the
routing layer that would point them at the references does not exist. End to end has never
run, so the performance regression is an expectation, not a measurement.
Section 6 writes down the three eager imports that blocked the plan, the annotation trap that
looks like it should solve it and does not, and the comparison table showing the quantisation
chain is bitwise equal -- plus, plainly, what is still missing so the next step is not mistaken
for a finished port.
The bring-up plan for a card that cannot build TileLang is to run the model on
`tile_kernels.torch` -- pure-torch implementations of nearly every kernel the forward pass
calls -- and replace the hot ones with Triton afterwards. That plan was blocked at the
import: `import tile_kernels.torch` dragged in the whole compiled stack.
Three eager imports were responsible, and each is now lazy or optional:
* the top-level `__init__` re-exported every compiled kernel; it resolves them on access
* `quant/__init__` did the same for 13 modules, and `import quant.common` runs it first
* `quant/common.py` mixed host-side config dataclasses with `@T.macro` fragments; the
dataclasses now import without TileLang, the macros are behind a guard, and its three
dtype-handle members resolve TileLang when called rather than when imported
`config.py` also only needed TileLang for three constants, now literals.
One thing that looks like it should work and does not: a module-wide
`from __future__ import annotations`. It also stringifies the annotations inside `@T.macro`
bodies, which TileLang's macro parser evaluates itself, and it fails there. Only the four
host-side annotations are quoted, individually.
The patches are idempotent and reversible, and `verify_tile_kernels_patch.py` proves the
surface is unchanged where TileLang is installed -- the kernel modules' wildcard imports of
`quant.common` still see the macros, `from pkg import sub` still works, config knobs still
take effect.
The other half of the question, whether the references are numerically trustworthy, is
answered by `compare_with_cuda.py`: on the recipe the engine actually uses on Hopper,
per_token_cast, cast_back, swiglu_forward_and_per_token_cast and engram_hash are bitwise
identical, and apply_rotary differs by one bf16 ULP on one element in 65536. The quantisation
chain being bitwise equal is what makes a torch-first bring-up safe rather than approximate;
a cast that were merely close would change tokens downstream.
Not yet done, and recorded in the doc: MHC and the MoE gate still need their own comparison
before anyone assumes they match, and the routing layer that points deepinfer's 25 call sites
at these references does not exist yet.
I had been reporting the port as "16 kernels left". That was wrong in a way that matters, so
this corrects it with a measurement rather than an estimate.
The measurement: make `tilelang`, `tile_kernels` and `tvm_ffi` unimportable via a meta-path
finder installed from sitecustomize -- sitecustomize specifically, because the eight DP ranks
are spawned and would not inherit an import in the parent. Then try to import the engine.
It fails at import, before any kernel is reached: `hugging_face/weights.py` pulls in
`layers/linear.py`, whose module body has `import tile_kernels`.
The reason the kernel ports did not help is that `tilelang_ops` only ever covered 44 kernels.
The model forward calls `tile_kernels` directly -- 25 module-level imports across 13 files,
23 distinct APIs -- and `tile_kernels` is itself a TileLang package: its `__init__` line 1 is
`import tilelang`, with 166 references across 140 modules. Keeping tile_kernels keeps
TileLang.
The 23 APIs are recorded in three groups, because they cost very differently: six are
configuration knobs and string helpers; six are hot-path quantization and routing kernels
(`per_token_cast` appears in four modules); ten are MHC/Engram kernels specific to this model
and the largest piece of work.
There is one real shortcut, and it is worth knowing before anyone starts writing Triton:
`tile_kernels/torch/` holds pure-torch reference implementations of nearly all of them, with
zero TileLang references. It is currently unreachable only because the package's `__init__`
imports TileLang eagerly -- so making that lazy buys a correct-but-slow bring-up path, from
which the hot kernels can be replaced one at a time as profiling justifies it. `flag_gems`
cannot stand in: it is 903 ATen-level operator replacements and does not carry these fused
kernels.
An edit of mine dropped a `return launch` from `finalize_verification_tl`. The A/B case for
that kernel was not re-run, so nothing local objected; the failure surfaced as a TypeError
inside the sampler of an end-to-end run, once per rank. The harness now checks that every
`*_tl` symbol the selected backend exports is callable, which turns that class of mistake
into one line printed before any case executes.
Also records where the remaining sixteen kernels actually sit. None of them is
unconditionally on the path, but the distribution is lopsided and the document had it
backwards in one place: the three `candidate_*` kernels are not merely "off on SM90", they
are switched off *by* `PLATFORM.is_sm90` -- and this checkpoint sets
`candidate_source_layer_id=20` of 40 layers. On a card where `is_sm90` is False, that is
nineteen layers on every step, which is the opposite of optional. Their harder blocker is
noted alongside: that path also needs `deep_gemm.fp8_fp4_sparse_mqa_logits`, which is
SM100-only, so porting those three alone would buy nothing.
The four kernels that no production path reaches are listed as such, so nobody spends an
afternoon on them.
`DEEPINFER_USE_NEXTN_DECODE` defaults to 5, so speculative decoding is on unless someone
turns it off, and these four ran on every decode step. On a host without TileLang the first
step failed; on the H20 they were the only kernels still answering from the other backend
during a normal run.
The interesting one is `remap_draft_topk`. Its TileLang form builds a vocabulary-windowed
lookup table in shared memory and reads the logprobs back out of it. Triton has no
equivalent primitive, so the table is gone and each target index is matched against the
draft indices directly: draft indices are unique per token id, so at most one candidate can
match and a max over the block is well defined. That is O(topk^2) instead of
O(vocab/pool_size * topk), which would be the wrong trade at a large `force_topk` -- but
`force_topk` defaults to 0, so this kernel is off the default path and correctness comes
first. Worth revisiting if a deployment turns it on.
The other three adapt rather than change shape:
* `get_accept_result`'s prefix scan becomes a bounded loop over the widest span in the
block; lanes with a shorter span hold their value on the extra iterations. MTP packs at
most num_steps+1 rows per request, so the bound stays small.
* `finalize_verification` reduces `batch_confidence_error` within the program and issues one
atomic instead of one per row, because a scalar pointer with a vector value is not valid
in this Triton version. Accumulating in a different order changes the last bits, so its
cases compare with a tolerance and say why.
* `get_fused_dflash_get_input` rounds its column loop up. `padded_kv_len` is a multiple of 64
but not necessarily of the tile width, and flooring leaves the tail of every row unwritten.
33 new A/B cases across the four, covering each constexpr variant the engine builds. The
full run is green and an end-to-end pass is token-for-token identical to the run before the
change, with the MTP fallback notices gone.
The dependency notes in section 4.4 were a list of things to do by hand; they now name the
script that does them, along with what a from-scratch run measured.
`h20_e2e_report.sh` sets the two settings a single-node H20 run always needs -- no GIN
because the container has no /dev/infiniband, and eager because the Engram gather cannot be
captured -- and appends Triton's own cudart to the loader path when it is present, since cu13
images otherwise fail to load any kernel. Requiring a caller to remember three exports to
reproduce a run is how a run stops being reproducible.
Getting the Triton backend to build and run in a vendor image took a dozen non-obvious
steps: a venv because setup.py insists on CONDA_PREFIX, CUDA headers borrowed from the
nvidia-cu13 wheels because the image ships only the runtime, NCCL 2.32.3 because DeepEP
uses ncclCommQueryProperties, libdw/libelf for DeepGEMM's bundled deep_jit, pciutils for
DeepSelect, and a tvm_ffi new enough for tilelang 0.1.15. All of it lived in this session
and in prose, which means the next person starts over.
`setup_flagtree_image.sh` records it in three stages -- env, build, verify -- and is
idempotent, so a failure resumes. It also explains why each workaround exists, since several
look arbitrary otherwise: FlagTree is a Triton fork and needs nothing special, but the C++
extensions around it do.
`copy_site_packages.py` copies distributions by their declared file lists rather than by
name, because pillow installs PIL and pillow.libs, pyzmq installs zmq. It defaults to
keeping the image's pillow: a copied one points at vendored libtiff builds the image does
not have, and PIL.Image then fails to import.
Verified by building into an empty venv rather than the one already set up here: env, build
and verify all pass from scratch, and an end-to-end run from that venv produces the same
eight-rank-agreeing output as before.
The raw log prints one block per rank, so the two things that decide whether a run passed
-- do all eight ranks agree on the token ids, and does the fallback list match the kernels
we know are unported -- are both buried in repetition. This extracts them: a count per
distinct token sequence (eight means every rank agreed), one continuation per prompt, the
kernels still served by the default backend, and whether a traceback occurred.
Splitting "hard failures" out from the stats logger's output matters on torch 2.11, where
`ProcessGroupNCCL.memory_stats` does not exist and the logger says so without the run being
affected.
Eight ranks, `--ep 8 --max_tokens 32`, both prompts, on the FlagTree image: exit 0, every
prompt's token-id sequence identical across all eight ranks, decoded text correct. The same
run on the reference image produces byte-identical token ids, so the port is not merely
"some output comes out" on the vendor image -- it agrees with the image we have been
validating against all along.
The fallbacks are the known P1/P2 leftovers and are identical on both images: MTP's four
kernels, the sampler's rejection-sampling variant, and the 20 optional-path kernels. One
benign difference is recorded too -- torch 2.11 has no `ProcessGroupNCCL.memory_stats`, so
a stats line reports itself unavailable.
Also writes down the dependency work an end-to-end run needs but the A/B harness does not:
tilelang 0.1.15 with tile_kernels 2.1.0 and tvm_ffi 0.1.11, plus the 19 packages from
requirements.txt. pip stalls on this host, so the practical route is copying site-packages
out of the reference image; pillow is the exception and must come from the image, because
11.3 sits beside a pillow.libs the image does not have.
Two corrections and one result.
FlagTree is not something Triton has to be ported to: `flagtree.__version__` reads
`triton.__version__`, and `triton/runtime/jit.py` imports `flagtree._flagprism`. The
distribution is named `tokenspeed-triton` but the import name is still `triton`, because it
is a fork. So "FlagTree cannot compile" was never true -- it compiled every kernel we have,
element-wise identically to triton 3.5.1 and 3.7.1. What failed was `deepinfer._C` and the
vendored CUDA extensions, none of which involve Triton at all.
Those are now built inside the FlagTree image (torch 2.11, cu13, gcc 13.3, nvcc 13.0):
`_C` loads with 34 symbols, deep_gemm 2.8.2, deep_ep 2.5.0 and deep_select 1.0.0 all
import, and the full A/B run is 125 cases with zero skips -- the sampler group used to be
skipped for want of `_C`. Comparing against the dumps from triton 3.7.1 is still exact.
Everything needed along the way is written down, because none of it is obvious: the venv
requirement in setup.py, `python` vs `python3` in build.sh, the missing cusparse header,
NCCL 2.32.3 for `ncclCommProperties`, pciutils for deep_select, and the pip download that
has to be a resumable curl to survive the 305 MB wheel.
Two artefacts and one finding.
`scripts/flagtree_triton_probe.py` asks whether the Triton an image ships can build the
features this port depends on, one feature at a time: `tl.dot`, `tl.gather` with a
data-dependent index, `tl.gather` as the RoPE pairing, `fp8e4nv` with a block scale, and an
int64 comparison (uint32 compares go signed in Triton, which the sampler depends on). Both
vendor images pass all of them and emit `.target sm_90a`.
`TODO.md` collects every open issue across both branches, with the two blockers an
end-to-end run in a vendor image actually hits: `deepinfer/_C` is built against one torch
ABI (`2.13.0+cu132`) so it fails to load on torch 2.11 and 2.9, and `deep_gemm`/`deep_ep`
are not on PyPI in the versions `vendor/` provides. Neither is a Triton problem, and saying
so plainly is the point of writing it down.
`h20_ops_ab.py` assumed the TileLang backend is importable, so on a vendor image it could
not start at all -- which is exactly where the Triton backend needs checking. Two modes fix
that:
* `--dump DIR` saves the selected backend's output per case.
* `--compare DIR` checks against dumps written elsewhere.
Together they answer "do triton 3.5.1 and 3.7.1 agree" without either host having TileLang.
Case builders that genuinely need a package the host lacks now report a skip instead of
aborting the run, so partial coverage is visible rather than fatal; the import of the
sampler's C++-backed inputs moved into the one builder that uses it for the same reason.
This is what showed the 19 ported kernels produce element-wise identical output across
triton 3.5.1, 3.6.0 and 3.7.1 -- 115 cases, no differences.
The point of the Triton backend is to run where TileLang cannot, but importing `deepinfer`
pulled the whole TileLang stack in regardless. Three separate leaks:
1. `ops.py` imported the platform default backend even when another one was selected, purely
to have fallbacks available. That default is TileLang, so it dragged in `tile_kernels`.
`DEEPINFER_OPS_STRICT=1` now skips it: no default import, no fallback, a missing kernel
raises. That is the right trade when the default is not installed, or when a silent
fallback would hide an incomplete port.
2. `tile_kernels` sat at module scope in `norm`, `cache_common`, `fused_transform` and
`load_paged_kv_cache`, but only the reference implementations and the unit tests use it.
The kernels never do. It now lives inside those functions.
3. `tilelang` itself: `compress` set a TileLang pass flag, and `norm` needed a TileLang dtype
handle to pass along a parameter the Triton factory discards. Both now degrade to a no-op
or a plain string when TileLang is absent.
Verified on a FlagTree image (torch 2.11, triton 3.6, no TileLang): `import` succeeds, 19
kernels build and run, and the outputs are element-wise identical to the same kernels built
by triton 3.5.1 and 3.7.1. Full A/B on the reference image is still green, so nothing
changed for the paths that do have TileLang.
`h20_ask.py` takes its prompts as arguments, and a login shell that is not UTF-8 mangles
them into `????` before python ever sees them. The result is a model answering a question
nobody asked, reported as a quality problem. The prompts now live in the script.
Also records, in the plan, the state of the port as measured today: what runs end to end,
the two environment variables it still needs, and the one real blocker found while
checking -- the Engram gather is not CUDA-graph capture safe, so the default
(attention-graph) configuration fails to warm up on eight ranks.
With MTP on (the default), the sampler is built with rejection sampling for the draft-verification
rows, which this backend does not implement. It asserted, so the whole MTP path failed with the
triton backend even though 16 of its 20 kernels were ported.
It now hands that variant to the default backend and says so once. That is a deliberate delegation
rather than a silent fallback: `fused_sample` is exported by this module, so the per-kernel
`_FallbackNotice` cannot cover it, and the note has to come from the factory.
Verified both ways on 8xH20, same tokens as the pure TileLang run either way:
MTP off 20/20 kernels Triton, no fallback notice at all
MTP on 16/20 kernels Triton, and exactly five notices -- the four unported MTP kernels plus
this delegation, with nothing unaccounted for
The remaining MTP work is `remap_draft_topk`, `get_accept_result`, `finalize_verification`,
`get_fused_dflash_get_input`, and this rejection variant.
compute_seqlens_k_zos, confidence_head_fused and dflash_input_store, all identical to the baseline:
15 cases across the three.
`compute_seqlens_k_zos` masks the remap lookup rather than clamping it: `-1` marks a fresh request,
and an unmasked `PREV_ACCEPTED + (-1)` would read the element before the buffer.
`confidence_head_fused` reads the weight as one flat vector and splits it at `n_embed`, so the
hidden and Markov halves are one accumulation instead of two passes. Without a Markov embedding the
second pointer is filled with the first tensor rather than left null, since the read is compiled out
but the argument still has to be a pointer. The scaled logit is clamped to +-60 before the sigmoid,
as the baseline does, so a small temperature cannot produce inf through `exp`.
`dflash_input_store` accumulates in fp32 even though both sides are bf16: the baseline sums into an
fp32 fragment, and a bf16 accumulator would round differently for hc_mult > 1.
One trap: `tl.next_power_of_2` inside a kernel body is not a compile-time constant, so the Markov
block size is resolved in the factory and passed as a constexpr.
The plan note still said the sampler was deliberately left on TileLang, which is no longer true. It
now carries the measured numbers instead: 100% token agreement on the engine's own configuration,
~90% on narrow cuts, log-probabilities within 1e-5 everywhere, and the reason the token cannot do
better -- the baseline's rejection limit steps at `total == 2^31`, which the fixed-point total sits
next to, so a few units of rounding error move to a different draw.
Plus the four traps that each looked like something else: signed arithmetic on uint32 (`/`, `%` and
`<=`), nested runtime `if`s losing the assigned value, a threshold taken from probabilities and
compared against logits, and renormalising by division instead of through the log domain.
`fused_sample` for the configuration the engine builds: 4096 columns after the force_topk cut,
sampling on, rejection off. With one tile the baseline's two-level search collapses -- the coarse
offset *is* the token id -- so this implements log-softmax, the top-k/top-p cut as a threshold plus
the tie rule, renormalisation, a fixed-point cumulative distribution and one Philox draw.
The draw matches cuRAND bit for bit (that is what the previous commit established). The rest does
not all match, and the split is measured rather than claimed:
engine's configuration (top_p=0.99, 2644 kept columns) tokens identical on every row
narrow cuts (top_k=128, top_p=0.5) ~90% of rows
Kept sets match column for column and log-probabilities agree to ~1e-5 in every case. What differs
is the token, and the cause is a discontinuity in the baseline rather than a defect here: the draw
is `d % total` with `total = sum(floor(p_i * 2^31))`, and the baseline's rejection limit
`(2^32-1) // total * total` steps from `2 * total` to `1 * total` exactly at `total == 2^31` -- which
`total` sits next to, because the probabilities sum to one. On the near side almost no draw is
rejected, on the far side about half are, and each rejection moves to the next draw. A few units of
rounding error in `total`, which is a reduction order that cannot match, therefore change the token
outright. Clamping `total` under 2^31 does not help: the baseline crosses in either direction too,
so it only relocates the disagreement (30 rows instead of 11).
Four bugs found on the way, all worth recording because each looked like something else:
- Triton evaluates `//`, `%` and `<=` on **uint32 operands with signed arithmetic**, so 4294967295
behaves as -1. That made `limit` come out as `2^32 - total`, flipping the rejection decision.
Fixed by doing the whole cumulative distribution in int64.
- Nested runtime `if`s produced a threshold that was not the one assigned inside the branch; the
cuts are now combined with `tl.where` instead.
- The threshold was taken from the probability array and compared against logits. The sort has to be
of the logits (equivalent, since the map is monotone -- the mass is recomputed from the sorted
logits rather than sorting twice).
- Renormalising by division rather than through the log domain, as the baseline does, moved the
token agreement from 76% to 91% before the remaining gap was understood.
Also gives the A/B harness a fourth case element: an agreement floor for outputs whose answer is a
random draw, so the rate is reported rather than compared for equality -- a comparison that can pass
without testing anything would be worse than none.
The last blocker for `fused_sample` was the random stream: `T.rng_init`/`T.rng_rand` lower to
`curand_init`/`curand` with `curandStatePhilox4_32_10_t`, and the seeding is not documented. Its
mixing function is public but `curand_init` folds the seed, the subsequence and the offset into the
counter in a way that has to be measured rather than derived.
Measured (`scripts/h20_rng_dump.py` dumps the real stream), the layout is:
key = (seed & 0xFFFFFFFF, seed >> 32)
ctr = (offset & 0xFFFFFFFF, offset >> 32, subsequence & 0xFFFFFFFF, subsequence >> 32)
with draws taking the block's four words in order before `ctr.x` advances. The subsequence sits in
the third and fourth counter words, not the first -- a cuRAND subsequence spans 2^64 blocks. That
detail is invisible at subsequence 0 and wrong everywhere else, which is why the check covers both.
`scripts/h20_philox_check.py` compares our kernel against the dumped reference over five seeds
(including 2^64-1) and four subsequences (0, 1, 2048, 4097 -- the last two are the actual
`sample_positions` from a run): 20/20 streams bit-identical.
One trap worth noting: the multipliers and the second Weyl constant all exceed int32, so referencing
them as Python ints inside the kernel silently wraps them negative. They are materialised as uint32
tiles instead.
This does not yet change any output -- `fused_sample` is still the TileLang kernel. It removes the
reason it had to be, which is the point.
The two runs agree token for token on the reference chat-prefix input, but a plain-text prompt
diverges after two tokens: both backends are internally consistent across all eight ranks, so the
difference is a different draw from a slightly different CDF, not nondeterminism.
That is the quantisation floor from P4 showing up through sampling rather than through argmax -- and
on top of it, a kernel written in a different language has its own reduction order (the norm A/B
shows 10 of 655360 bf16 elements one ulp apart, with fp8 bytes and scales identical). At T=1.0 that
ulp is enough to flip a coin.
So the honest statement is a table, not a claim: identical tokens on the chat-prefix input, same
first token but a different branch on a plain-text prompt, and matching the draw itself would
require both an identical reduction tree and the same PRNG stream.
The import-time line only counts kernels the selected backend does not export, which says nothing
about whether a running configuration reaches them. `_FallbackNotice` wraps each inherited kernel
and prints once, on first call, so an end-to-end run states the answer itself. On 8xH20 with
speculative decoding off that answer is a single line:
[ops] kernel 'fused_sample_tl' used from 'cuda': 'triton' does not implement it
Eighteen of the nineteen kernels on that path are Triton and produce exactly the same tokens as the
pure TileLang run. The nineteenth is left there on purpose.
A probe of the engine's own sampling parameters (`DEEPINFER_PROBE_SAMPLING=1`, which stages an
async D2H copy and reports at exit so it does not synchronise inside the forward) settles what
`fused_sample` is being asked to do:
temperatures: [1.0, 1.0] top_ps: [0.99, 0.99] top_ks: [-1, -1]
That is a real nucleus sample over 4096 candidates, not an argmax, so the draw matters. Two things
stand in the way of reproducing it bit for bit, both internal to TileLang: `T.rng_init`/`T.rng_rand`
are a Philox stream implemented in its C++ codegen, so a Triton kernel draws a different stream and
every sampled token at temperature > 0 would change; and the baseline pins the layout of its LSE
accumulator precisely to fix the summation order, after which probabilities are quantised into a
2^31-scaled uint32 CDF where a one-ulp difference can move a column across a boundary.
Swapping it anyway would be a behaviour change reported as a kernel swap. The per-kernel fallback
exists for exactly this case, so the sampler stays on TileLang and the rest of the path does not.
`triton/__init__.py` records this so the gap is not mistaken for an oversight, along with the order
in which porting it later would have to proceed: a Philox that matches the stream exactly, then the
radix filter, then the fixed-point CDF and the two-level search.
Four more kernels, and with them the non-speculative path no longer touches TileLang:
- `fused_norm_rope_hadamard_cast_store_kvcache` -- all the forms the engine builds: bf16 out,
bf16 out after dequantising (cast_back), quantised out with scales, and the main/index KV-cache
write with a bf16 stage of `norm`. 18 cases.
- `compress_and_update_state_{prefill,decode}` -- the ring-cache state machine, two launches each.
- `prepare_v4_decode_metadata` -- one program per sequence, every output including the
ratio-bucketed ones.
Details that were not obvious and are now pinned by cases:
RoPE pairs differ per layout: interleaved pairs `(2k, 2k+1)`, Neox pairs `(k, k + rotary/2)`. Both
reduce to "is this the first element of its pair" plus a partner offset, so the lane vector is
gathered once with `tl.gather` instead of two code paths. That also fixed a real bug: the first
version paired adjacent lanes for both layouts, which is only correct for interleaved.
`bf16_stage="norm"` means the post-normalisation *pre-RoPE* row -- the indexer consumes it. It has
to be stored before the rotation overwrites the row; storing it afterwards produced a wrong value
only when RoPE was enabled, which is exactly the case the engine uses.
Scales come in two storages for the same rounded power-of-two value: a ue8m0 byte for the main KV,
an fp32 for the index KV. Both are now covered, in the standalone output and in the cache write.
`compress` needed its tail-copy grid bounded by the *query token* count, not the compressed-row
count; and `_enable_fast_math` mutates `pass_configs` on the factory, a TileLang concept, so the
Triton factories expose an inert dict to keep that call working.
Several case bugs of the same shape were found and fixed along the way -- a sentinel that was never
required to be overwritten, a `total_q` that did not match the actual prefix sum, two tokens
mapped to one cache slot or one state slot so the winner depended on CTA order. Each of them made
the two backends agree for the wrong reason, which is worth more than the kernels themselves: a
comparison that can pass without testing anything is worse than no comparison.
The illegal memory access was not in `expand_sampling_params` at all, though that is where
`CUDA_LAUNCH_BLOCKING` pointed. It was the metadata kernel leaving its output unwritten.
`prepare_verification_metadata_tl` sized its second grid dimension from `num_steps`, i.e.
`max_MTP_value`. With speculative decoding off that value is 0, and a zero-sized grid dimension
launches nothing, so `row_to_req` -- a `torch.empty` buffer -- kept whatever the allocator handed
out. `expand_sampling_params` then used those values as request indices, and the reads went off
the end of the parameter arrays.
The kernel now launches one program per request and loops the tiles at runtime, mirroring the
TileLang version, so the row count comes from device memory and no host read is needed (which
would have been rejected as a synchronising operation anyway). `num_steps` is accepted and
ignored.
The A/B case for this kernel is also fixed, in two ways that matter:
- it used a fixed `total_q` while the request lengths were random, so it allocated more rows than
were ever written and the comparison was meaningless;
- it pre-filled the outputs with a sentinel but never required them to be overwritten, so both
backends could write nothing and still agree with each other. The case now asserts every row was
written, and covers `num_steps=0` plus a request longer than one tile.
With that, `DEEPINFER_OPS_BACKEND=triton` runs end to end and produces exactly the same tokens as
the pure TileLang run, `[17607, 2793, 3139, 666, 6127, 666, 320, 1]`, identical on all eight
ranks. Fifteen of the required kernels are Triton; the remainder still come from TileLang through
the fallback while they are ported.
Three fixes found by running the ported kernels inside the engine rather than only against the
A/B harness:
`map_lookup` took the row stride from the row width, which is wrong for the views the engine
passes (a slice of a wider buffer). Every kernel now takes its strides as arguments.
`expand_sampling_params` used `row_to_req` directly as an offset; padding rows carry a negative
request index, so that read out of bounds. It is part of the mask now, and `store_mtp_info` treats
a non-positive slot as "no write" for the same reason.
`prepare_verification_metadata` computed its grid from `seqlen_q.max().item()`, a device-to-host
read that the engine rejects inside the forward ("called a synchronizing CUDA operation"). The
grid now comes from the factory's `num_steps`.
Integration is not finished, and `H20-TRITON-PLAN.md` says so plainly: with
`DEEPINFER_OPS_BACKEND=triton` the engine runs the 15 ported kernels into the sampling phase and
then faults on an illegal access inside `expand_sampling_params_tl`, even though that kernel is
element-for-element identical to the baseline in its own A/B cases, and even though the negative
index, the reserved slot and uint64 seeds have all been ruled out. The note records the two
bisection steps worth trying next, the env-var tooling already in place, and the fact that one of
my own edits silently did nothing because the search string carried a trailing space.
Also records the two new Triton traps: an `other=` literal must match the destination dtype, and a
grid must never be derived from a device read.
store_kv_cache and load_paged_kv_cache, both byte-for-byte identical to the TileLang baseline on
the layouts SM90 uses: main/SWA KV (e4m3 values, ue8m0 scale bytes, 528 B per token) and index KV
(128 fp8 + one fp32 scale, 132 B per token). The store also covers fp32 scales, because the index
layout uses them. FP4 caches and the packed-ue8m0 index layout are refused instead of approximated.
Store cases compare the whole cache buffer, including the `-1` invalid-slot path and a block size
(134) that is not a power of two. Load cases cover dequantising to BF16, the index passthrough,
and holes expressed as `src_block_id == -1`.
Two traps worth recording:
The load kernel's `other=` default has to be a float literal. With an fp8 source pointer, `other=0`
fails to type-check with a message about casting an int32 tile to fp8e4nv, which points at the
wrong line.
The load cases originally filled the cache with random bytes. Random bytes decode as fp8 NaN and as
NaN scale factors, and NaN compares unequal to itself, so the harness reported mismatches that said
nothing about the kernel. The cases now build the cache with `store_kv_cache` and compare index
outputs as raw bytes, where a NaN bit pattern is just a byte pattern.
15 kernels and 62 cases agree.
norm, map_lookup, expand_sampling_params, prepare_verification_metadata, copy_block_tables.
Plus the measured parameterisation that made the port tractable: `DEEPINFER_TRACE_KERNEL_ARGS=1`
records each factory's construction arguments, which showed that e.g. `norm_tl` is only ever built
as one combination on this path (d=5120, bf16 source, fp8 out, per-128 fp32 scales, col-major,
round_sf, no residual) rather than the dozen the signature allows.
Three findings worth keeping:
`norm` has to round the fp32 normalised value back to bf16 *before* computing the per-block amax
and the scale. That is what makes its fp8 output bit-identical to quantising the bf16 result
separately, which is how the engine defines the reference. Without it the fp8 bytes scatter by an
ulp wherever a value lands on a rounding boundary -- which looked like a scale bug at first
(224 vs 240) and was not.
With `round_sf=True`, which is the production setting, fp8 output and every scale factor come out
byte-identical. With `round_sf=False` 240 of 327680 bytes differ by one ulp, because `amax/448`
depends on the order of a 5120-element reduction and two compilers will not agree on that order.
For the same reason the bf16 output differs on 10 of 655360 elements by one ulp. Both are
documented, and the A/B harness now picks its tolerance from the dtype instead of one global pair.
`map_lookup` resolves duplicate keys by taking the last match, not the maximum: the TileLang kernel
overwrites on every hit while its own Python reference uses `max`, and the two agree only when keys
are unique. Top-k indices always are, so the Triton version matches the kernel -- and the case
generator now produces unique keys, with a comment saying why.
13 kernels and 53 cases now agree with the baseline.
gather, gather_input_ids_zos, gather_engram_tokens_zos, expand_swa_index, kv_index_cat,
convert_kv_index, load_kv_index, store_mtp_info. 45 cases, every one bitwise identical to the
TileLang baseline through the shared interface function.
`scripts/h20_ops_ab.py` swaps `ops.<name>_tl` between the two backends and runs the public call
once per side, so the comparison exercises the interface layer's shape handling too, not just the
kernel. `scripts/ab_cases.py` holds the cases: real engine widths and dtypes, negative indices
for gather, -1 padding for the index kernels, page-boundary lengths, the slot-0 no-write path,
and the engram image-sentinel rule.
Five Triton semantics traps worth writing down, all documented in the plan:
- `tl.where` evaluates both arms, so `tl.where(i == 0, 0, LOAD(i - 1))` really reads `LOAD[-1]`;
that was an illegal memory access, fixed by clamping the index instead.
- A scalar pointer cannot carry a block-shaped mask; it has to be broadcast with
`+ tl.zeros((BLOCK,), tl.int32)` first.
- Two arms of one kernel cannot use different tile shapes, which is why `kv_index_cat` became two
launches rather than one fused kernel with a column-block branch.
- Triton's `data_ptr()` already includes a view's storage offset, so passing `out[:, w:w+k]` and
adding `w` again double-counts -- that showed up as 0x01010101 in the second half.
- Test-case capacity matters: `store_mtp_info` addresses `slot * nextn + step`, so the pool needs
`(max_slot + 1) * nextn` rows; too few rows let both backends write out of bounds and disagree.
Groundwork for replacing TileLang with Triton on the H20 path, plus the first ported kernel.
The scope is measured rather than guessed. `DEEPINFER_TRACE_KERNELS=<path>` proxies every `*_tl`
factory in the dispatcher and records its first use; one full text-generation run on 8xH20 says
exactly 24 kernels execute. `H20-TRITON-PLAN.md` lists them, in dependency order.
`DEEPINFER_OPS_BACKEND` selects the backend package (unset keeps the platform default, so nothing
changes by omission) and `DEEPINFER_OPS_FALLBACK` names individual kernels that must stay on the
default. Kernels the selected backend does not export at all are taken from the default with a
one-line notice, which lets a partial port run end to end the whole way through instead of only
at the end.
That coexistence is the reason for a parallel backend rather than an in-place rewrite: the
TileLang kernel is a validated baseline, so each Triton kernel can be diffed against it on real
shapes through the same public interface function, which is a much stronger check than comparing
against a fresh Python reference. `scripts/h20_ops_ab.py` does that swap and reports per case.
First kernel: `gather_tl`, all 9 cases bitwise identical (zeros filled for out-of-range indices,
int8/bfloat16 payloads, strided rows, int32/int64 indices, empty and single-row inputs).
Every comparison in this note feeds the reference's chat-template token ids, which is the
worst case for a quantization perturbation -- it is a continuation task with a single
high-probability token, so the margin at the contested position is only 1.64. Asking the same
model in plain text gets the first token right with no drift at all:
prompt: 问题：中国的首都是哪里？\n答案：
output: 北京。\n问题：中国的首都是哪里 tokens [6127, 876, 2056, ...]
Identical on all eight ranks. Without this note the document reads as though ordinary Q&A
carries the ** artefact, which would understate the port.
The note still ended at "next: P4" with an open checklist. Records the outcome instead: the three
decode candidates (KV cache, decode indexer, MTP) are all ruled out with measurements, and the
remaining flip is quantified -- noise std 1.63 against a reference margin of 1.638 at the one
position that differs, out of eight.
Also renumbers the sections, which had a stray unnumbered conclusion between 四 and 五.
Brings in the P4 investigation (per-step logits, KV round-trip and decode-indexer checks, the
recipe probe) along with the docs that record what each check ruled out.
The substantive bug is the debug_dump move: five platform-neutral files imported it through
h20_ops, which registers the SM90 import patches with override=True at module scope, so loading
any of them on SM100 replaced the real flash_mla and deep_select entry points.
The dump hooks live in five platform-neutral files -- `sampler.py`, `mlp.py`, `deepseek.py`,
`decoder_layer.py`, `model_runner.py` -- and all of them imported
`deepinfer.h20_ops.debug_dump`. That is wrong on every platform except SM90: `h20_ops/__init__.py`
calls `extend_import_patches(override=True, ...)` at module scope, so importing it from shared code
installs the SM90 import shims globally. On an SM100 machine that replaces
`flash_mla.decode_mla_mqa_sparse`, `flash_mla.prefill_attn_varlen_sparse` and
`deep_select.interface.topk` with the Hopper implementations, and on NPU it would do the same
before the Ascend shims get a chance to matter.
`debug_dump` itself has nothing platform-specific in it -- `os`, `pathlib`, `numpy`, `torch` --
so it moves to `deepinfer/debug_dump.py` and the five imports follow. Verified that importing it
no longer pulls `h20_ops` into `sys.modules`, and re-ran the end-to-end check on H20: same output
token for token, `[17607, 2793, 3139, 666, 6127, 666, 320, 1]`.
This bug is already in the open PR, so it needs to land before that PR is worth reviewing further.
P4.2 completes the decode comparison, and it closes the investigation rather than opening
another thread.
The one prerequisite was `DEEPINFER_USE_NEXTN_DECODE=0`. Its default is 5, so MTP speculative
decoding is on by default, which means several draft/verify rows per step and `logits[0]` no
longer addresses the position you think it does -- that is why an earlier attempt misread the
pos-15 vector as pos-14 and looked like a missing step. With MTP off, one row per step, clean
alignment. Disabling it is also the right comparison, since the reference `generate.py` has no
MTP at all.
Per position, 11-token prompt, 8 decode steps:
pos ours ref ours margin ref margin d margin
11 17607 17607 10.695 9.861 +0.835 (prefill)
12 2793 2793 8.448 10.328 -1.880
13 3139 3139 15.276 11.946 +3.330
14 666 6127 1.480 1.638 -0.157 <-- only flip
15 6127 6127 15.950 15.109 +0.842
16 666 666 20.118 17.874 +2.244
17 320 320 9.829 10.402 -0.573
18 1 1 12.262 11.091 +1.171
d margin: mean +0.711, std 1.628, max 3.330. Seven of eight positions are right. The single
flip sits where the reference's own margin is 1.638, against 9.9-17.9 everywhere else -- the
noise (std 1.63) and that margin are the same size, so it is a coin toss. The reference takes
`6127` (北京) first, we take `666` (`**`) first, and both go on to produce the same text:
`**北京**` is valid markdown and means the same thing.
So the remaining difference is per-128 quantization, which SM90 gets no choice about, landing on
a position where the model itself is undecided. Not something to fix: the only lever is the
lossless e8m0 per-32 recipe from P3.6, and DeepGEMM's `gran_k == 128` assertion forecloses it
(see `H20-DRIFT-ANALYSIS.md`).
Also rules out two more suspects. MTP is not the cause: with `USE_NEXTN_DECODE=0` the output is
byte-identical, `[17607, 2793, 3139, 666, 6127, 666, 320, 1]`. And no logits post-processing is
involved: at pos 14 our top-1 is 666 with 6127 second by 1.480, so 6127 cannot have been masked
-- it would not be the runner-up if it had been.
Records one methodological correction: the original criteria assumed prefill and decode should
agree, and treated a large rel_err between them as a decode defect. The reference's own
prefill-to-decode transition moves the margin by about 2x, because the KV cache goes through a
quantized round trip. The meaningful comparison is d margin against that position's margin, not
rel_err on its own. The old criteria stay in the document as process history, marked revised.
Two closed-loop checks, both clean, which finally retires the two standing suspects.
P4.3 KV cache. The repo's own store/load tests pass for the fp32-scale index layout SM90 uses
(6 each). Running every layout gives 48 passed / 12 failed, and every failure is a ue8m0 index
layout -- `FP4_INDEX_KV_FORMAT` or `MXFP8_INDEX_KV_FORMAT`, neither of which SM90 selects. Those
failures are in the test's reference, not the kernel: it reinterprets ue8m0 bytes as
`PLATFORM.fp8_sf_type`, which is float32 here, so `index_copy_` hits a dtype mismatch.
`scripts/h20_kv_roundtrip.py` then closes the loop without any reference at all -- write known
BF16 through the paged cache, read it back, compare. Both SM90 layouts, block 32 and 64, 1/233/
1029 tokens:
V41_FORMAT (main/swa KV) rel 0.0266-0.0295 cos 0.9996
FP8_FLOAT32_SF_INDEX_KV_FORMAT (132B) rel 0.0266-0.0327 cos 0.9996
That is the fp8 floor, the same order as the 2.7e-2 on weights, so the cache is not where the
decode error comes from and the 132-byte layout question from P3.9 is closed.
One trap: the first version of the script compared raw fp8 values against BF16 and reported
rel=126. Index layouts are deliberately *not* dequantized by `load_paged_kv_cache` -- the
consumer is a kernel that wants bytes plus scales. Multiplying by the scale gives 0.027. The
127x was the test, not the implementation.
P4.4 the decode indexer. `_decode_p_sum_sm90` reimplements DeepGEMM's paged MQA logits in torch
because Hopper's kernel requires 64-token pages; nothing else exercised it. Compared against a
float64 reference that shares no code -- explicit per-token offsets instead of strided views --
it agrees to rel~1e-7 / cos 1.0 across page sizes 32 and 64, boundary-crossing lengths,
multi-page contexts and 32 heads. The `q_sf[..., :1]` dequantization is also sound: head_dim is
128, equal to the granularity, so one scale per (row, head) is correct.
What this changes. `logits_prompt` from the max_tokens=1 run puts the prefill error at cos
0.9846 with an absolute error of 3.30 on the winning token, against a reference margin of 3.112:
the margin is thin because the error is the same size as the margin, not by coincidence. The
reference's own prefill-to-decode transition also costs about 2x margin (3.112 -> 1.638) purely
from the quantized KV round trip. Our 0.256 halved would still be positive, yet it flips, so
roughly 0.5-1.0 of logit-level error in decode remains unexplained -- now narrowed to attention
over the cached main KV and the SWA-window / compressed-position mapping.
Also corrects an earlier overstatement: "the difference is introduced by decode" does not mean
decode has a discrete bug, and P4.3/P4.4 show its two main suspects do not.
One more dump trap worth recording: with max_tokens > 1, `logits_prompt` is contaminated too.
mysteps6 (max_tokens=7) gives cos 0.68 and a 16.5 top-1 error for the same input where mylogits3
(max_tokens=1) gives 0.9846 and 3.30. Prefill logits are only trustworthy from a
max_tokens=1 run.
Adds `record_step` so every decode step is kept instead of just the first, a `--force_ids` mode
to `ref_next_token.py` that mirrors `generate.py`'s incremental loop, and four fixes to the dump
path without which the data was unusable:
- filenames carry `_rank<R>_pid<P>`: with data parallelism every engine is its own
single-process group and reports RANK=0, so eight processes wrote the same file and the last
one won;
- decode reads row 0, not the last row: the batch is padded and the padding rows hold logits for
placeholder tokens, so the last row was a constant `top1=69, margin=0.0528`;
- only the driver rank records: a follower only joins the embedding all-gather, so its logits are
placeholders that got recorded dozens of times and buried the real steps;
- `DEEPINFER_DUMP_TOKENS` gating plus `profiling()` now cover both profile paths.
The conclusion itself comes from a black-box four-way comparison rather than the dumps, because
per-step dumps still do not line up with the emitted token sequence (`max_tokens=7` produced 75
sampler calls). At the same 14-token context:
reference prefill 6127 (北京) correct
reference decode 6127 (北京) correct
our prefill 6127 (北京) correct
our decode 666 (**) wrong
So prefill is right and the remaining flip is introduced by decode alone. Note the reference's
own prefill and decode margins differ (3.112 vs 1.638), so the two paths are not expected to be
identical -- KV cache goes through a quantized round trip -- but they should not disagree on the
direction.
Remaining suspects are therefore narrowed to decode-only state: the KV cache round trip (P4.3)
and the indexer scored in torch versus the TileLang prefill kernel (P4.4), which have never
been diffed against each other.
Comparing final logits was impossible before: `ref_layer_dump.py` read `logits[0, -1]`, but the
head already collapsed the sequence axis (`x = x[:, -1]`), so logits are `[b, vocab]` and that
index silently stored a 0-dimensional scalar. Every earlier dump had an unusable
`final_row_logits`.
With that fixed, the comparison reframes the remaining drift. On the same 14 tokens:
prefill top-1 = 6127 (北京) margin(6127 - 666) = +0.256
reference top-1 = 6127 (北京) margin = +3.112
Our prefill picks the right token; the error only compresses the margin, by ~12x. But the
end-to-end run, which reaches that context via decode steps, picks 666 (`**`) instead. Same
context, different path, different answer -- so the remaining drift is in **decode**, not in the
per-128 quantization. The quantization floor is still 2.7e-2 and still irreducible
(`H20-DRIFT-ANALYSIS.md` has the DeepGEMM assertion that forces `gran_k == 128` on SM90), but it
is not what flips the token.
Two traps had to be handled to get numbers that mean anything:
- `execute_model_profile_run` runs a full forward + sampler on synthetic tokens. The layer dumps
were filtered by their prompt-length gate; logits are not, because their row count is the
number of *selected* tokens (1). The first attempt captured that step and produced cos = 0.34.
`debug_dump.profiling()` now suspends capture for the whole profile step.
- With `max_tokens=1` both a prompt and a decode step run, and only the prompt one is comparable
with a prefill-only reference. They are dumped as `logits_prompt` / `logits_decode`.
Also re-runs the layer comparison now that the expert-scale fix has landed: layer 0 went from
1.27e-1 to 2.88e-2 and layer 39 sits at 2.00e-1, so the error curve flattened by an order of
magnitude.
`h20_quant_check.py` cast to `torch.float4_e2m1fn_x2` to cross-check the code table, nibble order
and quantize/dequantize round trip. That cast kernel exists neither on SM90 nor on the CPU -- it
raises a device-side assert on H20 and `copy_kernel not implemented` on CPU -- so the check could
not run at all on the hardware it is meant for.
The E2M1 table and the unpack order now come from `inference/convert.py` (inlined, with the
source noted), and quantization picks the nearest table entry instead of casting. Still pure
integer/table work, so it runs anywhere and the round trip reports the same 2.77e-2 that the
expert-weight diff attributes to the per-32 -> per-128 requantization.
No behaviour change. `CONTRIBUTING.md` requires `ruff check --no-cache` and
`ruff format --check --no-cache` to pass, and the SM90 branch predates that requirement:
72 findings, all in the files added here.
Most of it is mechanical (`ruff format`), plus:
- annotations for the helper signatures (`torch.Tensor` / `int` / `Generator` / return types);
hooks are typed as `Callable[[nn.Module, tuple, object], None]` rather than left bare;
- two `if`/`else` assignments become ternaries (SIM108), and the index read in
`h20_real_requant_check.py` uses a context manager (SIM115).
Also corrects the `sparse_mla_fwd` docstring, which still described the old
`0 <= idx <= q_position` masking contract. Passing `is_causal=True` there is what produced the
all-NaN layer 0: indices address the assembled `[swa cache | current swa | compressed]` buffer,
so the row number is not a valid causal bound and the engine must pass False.
Consolidates the whole bring-up into one document so the next person does not have to
rediscover it: timeline by phase, every trap with symptom/root cause/fix, the full change
inventory, a reproduction guide, and an explicit TODO for the remaining numerical drift.
The three traps worth reading first are methodological, not code bugs, because each of them
made the comparison numbers meaningless rather than wrong:
- the reference wraps prompts in a chat template, so identical *text* is not identical *input*;
- `decoder_layer` returns the sub-layer output, not the residual stream, so comparing without
`hc_post` showed cos=0.02 at every layer;
- the engine runs a dummy profile forward before allocating the KV cache, so an ungated dump
captures the wrong token count.
Also records the drift TODO with measured numbers: the expert weights are at 2.7e-2 rel_err,
which is the per-32 -> per-128 requantization floor rather than a defect, but the reference
converter is lossless so P3.6 could still recover it if SM90 DeepGEMM accepts an e8m0 per-32
expert scale. P4/P5 in H20-DEBUG-PLAN.md are now cross-referenced to the journal.
`requant_expert_fp4` expanded the UE8M0 scale with `_expand_blocks`, which repeats *both*
axes. That is right for dense weights, whose scales are `(n_out/32, k/32)`, but expert scales
are already per-output-channel: `(n_out, k/32)`. Expanding both axes therefore treated
`(2304, 160)` as `(72, 5)` and handed every one of the 2304 output channels the scales of the
first 72 rows, scrambling each expert's weights.
This is what the layer comparison was pointing at: `shared_expert` matched the reference
(cos 0.9987) because it goes through `requant_dense_fp8`, which uses `_expand_blocks`
correctly, while `combined` -- routed experts plus shared -- was off by 3.6e-1.
Verified against the authoritative implementation, `inference/convert.py`'s
`cast_e2m1fn_to_e4m3fn`, on layer 0 / expert 370:
w1 cos 0.8936 -> 1.0000 rel_err 5.35e-1 -> 2.84e-2
w3 cos 0.8960 -> 0.9997 rel_err 5.26e-1 -> 2.69e-2
w2 cos 0.8964 -> 0.9996 rel_err 5.20e-1 -> 2.65e-2
The residual ~2.7% is the per-32 -> per-128 requantization loss, not a defect; the reference
is lossless because it folds the relative per-32 offset into the FP8 values.
End to end on the 11-token chat prompt:
before: [17607, 41175, 389, 666] "中国的首都是**"
after: [17607, 2793, 3139, 666, 6127, 666, 320, 1] "中国的首都是**北京**。"
ref: [17607, 2793, 3139, 6127, 320] "中国的首都是北京。"
The answer is now correct and the first three tokens match the reference exactly. The remaining
difference is markdown emphasis (`**`), so some numerical drift still moves the top-1 on the
occasional position.
Also adds two checks that would have caught this immediately:
`scripts/h20_expert_weight_check.py` diffs our requantization against DeepSeek's own converter,
and `scripts/h20_quant_check.py` pins the E2M1 table, nibble order and UE8M0 bias. The nibble
order and table turn out to match the reference; only the scale broadcast was wrong.
Narrows the SM90 numerical divergence to layer 0's FFN. The sparse MLA attention
kernel is not the problem: at layer 0 its output matches the reference with
cos=0.99897 / rel_err=4.7e-2, while the full layer output is already at 12.7% and the
error compounds to cos=0.865 by layer 11, which finally flips the top-1 token.
Tooling
- `h20_ops/debug_dump.py`: opt-in capture (`DEEPINFER_DUMP_LAYERS=<dir>`), last
position only, reduced over `hc_mult`, written to `.npz` at exit.
`DEEPINFER_DUMP_TOKENS=<n>` gates on the real prompt length, because the engine
runs a dummy profile forward with a different token count before it allocates the
KV cache and that dump would otherwise be mistaken for the real one.
- `scripts/ref_layer_dump.py`: the same anchors from the reference model via forward
hooks (block output, attention sub-layer output, embedding), plus the final
logits and greedy token.
- `scripts/compare_layers.py`: cosine similarity, relative L2 error and norm ratio
per layer, reporting the first layer above a tolerance.
Two traps this had to fix before the numbers meant anything
- `decoder_layer` returns the *sub-layer* output; the residual stream lives in
`residual` (`HCLazyPostData`, hc_mult=4). Comparing without merging them through
`hc_post` produced a bogus cos of 0.02 at every layer.
- `DEEPINFER_DEBUG_NAN` and the dump both need a host copy inside the forward, so
both take `allow_cuda_sync()` rather than fighting the engine's sync guard.
Also documents the reproduction case and the findings in H20-DEBUG-PLAN.md.
DeepInfer targets Blackwell: the vendored FlashMLA fork only ships SM100 kernels
(`CUDA_ARCHS = ["sm_100a", "sm_103a"]`, every entry point asserts `arch.is_sm100f()`),
and several DeepGEMM attention APIs plus the whole quantization recipe assume
FP4 tensor cores, packed UE8M0 scales and per-32 blocks. None of that exists on
Hopper, so this branch adds an SM90 backend and gets an 8x H20 node to produce
tokens end to end.
Scope and status: the pipeline runs (prefill, decode, sampling) and matches the
reference on the first token, but numerics still diverge from position ~12 onward,
so the output is not yet correct. H20-DEBUG-PLAN.md tracks the remaining work with
the exact reproduction case; see H20-SM90-FINDINGS.md for the analysis.
Quantization
- `Platform.sf_pack_factor` separates two meanings that `fp8_sf_type` conflated:
the GEMM scale type (FP32 on SM90) and the number of scales packed per element
(1, instead of 4 UE8M0 exponents per int32). KV caches keep the ue8m0 byte layout.
- Requantize FP4 experts and ue8m0 dense weights to FP8 E4M3 with FP32 per-128
scales while loading (`hugging_face/weights.py`, `h20_ops/quant.py`), so the
runtime sees a plain FP8 checkpoint: `mxfp4` is disabled at runtime, expert
weights are repacked as 2D, and `use_fp4_sparse_kv` is off since the FP4 cache
kernels do not exist on Hopper.
- GEMM recipes switch to `(1, 128, 128)` with `disable_ue8m0_cast=True`, which is
the only shape SM90 DeepGEMM accepts.
Attention
- `h20_ops/sparse_mla_sm90.py` adapts TileLang's `examples/deepseek_v32` sparse MLA
kernel (which is Hopper-targeted) to the V4.1 layout: RoPE lives inside the
512-wide latent, so the V3.2 `dim`/`tail` split is removed and QK/PV are both 512.
Measured on H20: 104 TFLOP/s, bf16-noise accuracy. It takes no causal mask,
because the indices are already causal -- deriving causality from the row number
drops every index once the KV buffer is offset, which produced an all-NaN layer 0.
- `h20_ops/attn_sm90.py` re-implements `prefill_attn_varlen_sparse` and
`decode_mla_mqa_sparse` with the same signatures as `attention_backend.py`, and
registers them through `patch_imports`, so call sites are untouched and the
FlashMLA-only module is never imported.
- Softmax normalisation clamps `sumexp`, so a padded row with no valid index yields
zeros instead of 0/0 and cannot poison the whole hidden state.
Engram
- `h20_ops/engram_local.py` replaces `deep_ep.EngramBuffer.fetch`, which needs
NCCL GIN (RDMA) and is unavailable on a single node, with a distributed gather
over NVLink built from two `all_to_all`s. Row ownership follows the C++ layout
(`owner = idx // num_local`). Verified against a full reference table on 8 ranks.
Other Blackwell assumptions removed
- DeepSelect's top-k kernel faults at launch on SM90, so `h20_ops/deep_select_sm90.py`
reproduces its contract with `torch.topk` (end truncation, OOB fill, in-place
output views). The shared-layer candidate-block optimisation is disabled because
`fp8_fp4_sparse_mqa_logits` has no SM90 implementation, and the dense MQA-logits
path is used instead; its fp32 output is narrowed for the block-max kernel.
- Decode-time index scoring uses `_decode_p_sum_sm90`: Hopper's paged kernel
hard-requires 64-token pages while the cache uses 128, and the pages are strided
views that cannot be split without a copy.
- `mega_mhc` / `mega_moe`, the FlashAttention-4 CuTe DSL kernels and the fused
norm+rope+cast kernel are all SM100-only and take a Hopper fallback.
- `allow_cuda_sync()` is used for the few places that genuinely need a host read
(all-to-all sizing, the ViT fallback, the NaN diagnostic) instead of disabling
the sync guard globally.
Vendored submodules (not part of this commit; submodule gitlinks are unchanged)
- `vendor/DeepSelect/setup.py`: re-enable the `sm_90a` gencode that is commented
out, otherwise top-k fails with "no kernel image".
Scripts: `h20_attn_check.py`, `h20_engram_check.py`, `h20_topk_check.py`,
`h20_smoke_check.py`, `h20_real_requant_check.py` are unit checks that run on 8x
H20; `h20_ask.py` is a one-prompt harness and `ref_next_token.py` produces the
official reference next-token for arbitrary ids, so both sides can be compared on
identical input.
