loopyt 21d051c4ad debug: per-step decode dumps, and prove the flip is decode-only
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.
2026-10-21 19:07:49 +08:00
..
tilelang_ops feat: add SM90 (Hopper) support for DeepSeek-V4.1 inference 2026-10-21 15:18:56 +08:00
__init__.py initial commit 2026-10-20 12:09:42 +08:00
attention_backend.py initial commit 2026-10-20 12:09:42 +08:00
layernorm.py initial commit 2026-10-20 12:09:42 +08:00
linear.py feat: add SM90 (Hopper) support for DeepSeek-V4.1 inference 2026-10-21 15:18:56 +08:00
mtp_info_pool.py initial commit 2026-10-20 12:09:42 +08:00
rotary_embedding.py initial commit 2026-10-20 12:09:42 +08:00
sampler_utils.py initial commit 2026-10-20 12:09:42 +08:00
sampler.py debug: per-step decode dumps, and prove the flip is decode-only 2026-10-21 19:07:49 +08:00
vocab_parallel_embedding.py initial commit 2026-10-20 12:09:42 +08:00