MI355 GPT-OSS-120B: fix six deadlocks and two silent MFMA hazards; recover the latency cost; add perplexity harness - #6
Merged
sangeeta0201 merged 25 commits intoAug 5, 2026
Conversation
Attention chunks in the fused full-layer gang task are claimed by xcd_rank, but Phase 3 was nested inside the Phase 1 QKV guard (`xcd_rank < total_qkv_tiles_per_xcd`, =10 for GPT-OSS 120B). With NUM_KV_CHUNKS > 10 only 10 workers ever reached the chunk barrier, so the `(s_chunk_prev % NUM_KV_CHUNKS) == NUM_KV_CHUNKS-1` merge condition never fired and the megakernel deadlocked. Since the host selected 16 chunks for seq > 1024, every run above 1k hung. Worker supply is set by the rank guards, not by grid_dim. Un-nest Phase 3/4/5 from the QKV guard so all workers_per_xcd (30) workers can serve chunks. Workers that skip Phase 1 also skip its epoch wait, so they get a catch-up wait before reading K/V, and qkv_epoch_expected is now snapshotted before any worker can bump the epoch. With the barrier fixed, scale chunks with sequence length instead of the fixed 8/16 split: per-chunk attention work is proportional to seqlen/NUM_KV_CHUNKS, so holding it constant keeps decode latency flat. Empty chunks already stamp LSE=-inf, so over-provisioning at short seqlen is safe. Add a [CFG] line and a guard that rejects a chunk count above workers-per-XCD. Wall-clock ms/iter (serial, GPU 0), before -> after: 1024 2.295 -> 2.295 2048 hang -> 2.334 4096 hang -> 2.464 8192 hang -> 2.582 16384 hang -> 2.722 32768 hang -> 3.040 49152 hang -> 3.401 65536 hang -> 3.714 1.6x growth across a 64x sequence increase, with coherent output throughout. Note: the `Decode: avg ms/iter` diagnostic under-reports at long sequences -- FWDPASS_LOG_MAX (persistent_kernel.cuh) caps the device timing buffer at 8192 entries, so beyond 8k it averages only the first 8117 iterations. The figures above are wall-clock, covering every iteration. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
… workers The fused full-layer gang task (type 216) deadlocked intermittently on GPT-OSS 120B / MI350 -- 4 hangs in 7 runs at seq 1024, chunks=8 -- independent of chunk count and sequence length. Nothing printed after "[HOST_DBG] Starting poll loop". Three barrier targets are snapshotted at task entry as "read_counter + 1". That is only valid if the worker reads before this layer's producer bumps the counter. qkv_epoch_expected was safe: its bump is gated on every participant arriving. attn_release_expected and routing_expected were not -- their producers are the merge worker and the TopK completer, not the consumer set, so nothing ordered a consumer's read against them. All 36 layers run inside one task (the ml loop in persistent_kernel.cuh) with only a per-block __syncthreads between layers, so workers skew freely, bounded only by barriers they mutually arrive at. Phase 2's participant count was max(total_qkv_tiles_per_xcd, NUM_KV_CHUNKS) -- 10 or 16 of 30. A worker above that rank arrived at nothing before Phase 7b; it only waited. It could therefore fall a full layer behind, read an already-bumped counter, and wait for a bump that layer never produces, while the layer that would produce it waited on that worker's MoE tiles. Widen Phase 2 to all workers_per_xcd. It becomes the single point every worker on the XCD passes through, pinning them all to the same layer before either producer can run, which makes both "+1" snapshots well-defined. Safe because the dispatcher gives this task exactly workers_per_xcd tiles per XCD (gang_task_tiles_per_xcd = params[28]; dispatch_count = min(n_tile_count, workers_per_xcd)), one per worker, so every rank arrives exactly once. Workers with no QKV tile arrive immediately and add no latency. Validation, strictly serial on one GPU: seq 1024 chunks=8: 4 hangs / 7 runs -> 0 hangs / 10 runs seq 4096 chunks=16: 0 hangs / 4 runs Latency unchanged: 2.34 ms at 1k (was 2.32), 2.47 at 4k, 2.56 at 8k. Not fixed here: a separate pre-existing "Memory access fault ... on address (nil)" crash (~1 in 6), distinguishable by the process aborting during the first decode iteration rather than hanging silently. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Two defects found while auditing the fused-layer path: 1. layer_idx was stored as 0 into the LDS slot on every task entry and never advanced, so the MoE W13->W2 barrier's release value (layer_idx + 1) was permanently 1. Its d_barrier is monotonic and never reset, so once layer 0 wrote 1 the barrier was already satisfied for every subsequent layer: W2 workers stopped waiting for their own layer's W13 and could read swiglu_out before it was written. Tile ordering usually hid this, but nothing enforced it. Both live callers now publish a counter that actually advances once per layer and is never reset: qkv_epoch_expected in the full-layer fused task (216), and the routing epoch in the oproj+topk+moe fused task (217), which never initialized the slot at all. 2. FWDPASS_LOG_MAX = 8192 silently truncated the device-clock timing ring. Because per-iteration latency grows with sequence length, a 16k run reported the average of its cheapest 8k iterations. Added untruncated aggregates (total_ns / total_iters / dropped) printed as [FWD_PASS_TOTAL], and a host-side NOTE when samples were dropped. Verified at 16384: iters=16382 dropped=8191, all-iteration average 2.663 ms vs 2.280 ms reported by the truncated ring. 1024 unchanged (5/5 clean, 2.326-2.373 ms). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The multi-layer fused path reloads task_desc->input_ptrs/output_ptrs from ml_input_table/ml_output_table for every layer past the first, and dereferences every slot it reloads. A null slot there surfaces only as "Memory access fault ... on address (nil)", which aborts the process with no usable wavefront state -- device printf output is lost on abort, and the GPU coredump ROCm writes carries no register state. Check the tables on the host instead, where a bad slot can be named by (layer, xcd, index). Runs at graph-build time only. Current result on this model: 0 null inputs, 0 null outputs, so this rules the tables out as the source of the nil-address fault rather than fixing it. That fault is still open. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The nil-address memory fault aborts the process, and every ordinary
diagnostic channel is unavailable on that path:
- rocgdb perturbs timing enough to hide the fault entirely (7/7 clean
against a ~1-in-6 base rate), so running under the debugger is not
an option.
- Device printf output is lost on abort. Faulting logs contain no
[PC_INIT] or [WORKER_XCD] lines at all.
- The GPU coredump ROCm writes (~34 GB) carries no wavefront state:
no active dispatches, all registers <unavailable>.
This adds an opt-in tripwire (MPK_NIL_TRIPWIRE=1) that records, per
worker, the layer / phase / tile / base pointer it last reached. The
buffer is pinned host memory mapped for device access, so writes land
in host RAM as they happen and survive the abort.
Two details worth keeping:
- Only plain stores to each worker's own slots. A first version
scanned all 35 pointers per layer and cost 4.5x (10.67 vs 2.33 ms),
which would have hidden the race exactly like rocgdb. The cheap
version measures 2.47 ms at 8k, matching baseline.
- A 100ms disk snapshot thread, not just a SIGABRT handler. The ROCm
runtime installs its own SIGABRT disposition after ours, so on a
real fault our handler never runs (observed). The file survives
regardless of who wins that race.
The fault has been reproduced with the tripwire armed, so it does not
suppress the bug. Still unfixed; this is the instrument, not the fix.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The existing tripwire only resolved to "inside _execute_gang_task", which
is an entire fused layer. At that resolution a worker that faulted is
indistinguishable from one merely parked at a barrier, so the breadcrumbs
could not say where the nil address came from.
Widen the per-worker slot count 4 -> 8 and add:
- MPK_TW_SUB markers through the 8 phases of the fused layer kernel
(QKV, barriers, attention, merge, cross-XCD, O-proj/TopK, MoE, exit).
- Markers inside the MoE kernel, which is where the fault landed. These
hang off the existing MOE_DBG_SUBPHASE no-op hook, plus two new
pre-decode markers (1000/1001) covering the routing-mask read itself.
aux packs the decoded tile identity -- global_tile, expert_idx,
expert_id, num_activated_experts, is_w2 -- because every address the
kernel computes derives from those.
- The decode iteration (pc_iter) at layer entry. The host only sees that
it died within a second of launch, which at ~2.4ms/iter is anywhere in
the first few hundred; the first catch was at iteration 7903 of 8192.
g_tw_dev is published by the worker loop because task kernels take pointer
arrays rather than the RuntimeConfig and cannot reach config.tripwire.
Every block stores the same value, so the race is benign.
Cost is ~2% (2.399 vs ~2.35 ms/iter at 1024). That margin matters: rocgdb
perturbs this fault enough to hide it entirely (7/7 clean vs ~1-in-6), and
an earlier tripwire that scanned all 35 pointers per layer cost 4.5x and
would have done the same.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The TopK worker publishes routing_ready via st_wt (write-through, bypassing L2) to wake MoE workers on all 8 XCDs. Immediately before, all 256 threads of that block wrote active_expert_ids and routing_indices with ordinary st_wt stores. Only a __syncthreads separated the two. That orders the stores within the block; it says nothing about when they become visible to another XCD's L2. So the release flag can reach HBM ahead of the routing data it advertises, and a remote MoE worker can pass the Phase 7b barrier and read a stale active_expert_ids / routing_indices. The contract was already documented in two places and simply not implemented: the comment here refers to a threadfence_gpu that was never in this function, and gang_oproj_topk_moe_fused_mi300.cuh states "TopK worker wrote per-XCD flags via st_wt after threadfence_gpu". Impact is silent numerical corruption, not a crash. A stale read yields wrong values that are all still in range: the count at [NUM_EXPERTS] is the compile-time constant k=4 on every layer, expert ids stay in [0,128), and route_val stays in [0,4]. So a token can be routed through a previous layer's expert -- wrong output, no fault. That last point is worth being explicit about, because this fix was found while hunting the nil-address fault and it does NOT explain it. Every address the MoE kernel derives from the routing data stays in bounds even when the data is stale. The nil-address fault remains unexplained. Costs one fence per layer: 2.441 vs 2.399 ms/iter at 1024 (~1.7%). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The ml loop rewrites task_desc->input_ptrs / output_ptrs in place for each layer, but copied only the first 24 inputs and 11 outputs -- sized for the plain fused-layer kernel. The LM-head variant (TASK_GANG_FULL_LAYER_WITH_LMHEAD_FUSED_MI300) reads input_ptrs[24..27] and output_ptrs[12]: the final norm weight, the LM-head MXFP4 weight, its bias, and the argmax output slot. With FUSE_TAIL=1 those five slots would never be refreshed past layer 0, so every later layer would dereference layer 0's pointers -- and TaskDesc is reused shared memory, so a slot the plain variant never writes holds whatever the previous task left there. Latent in the default configuration: FUSE_TAIL defaults to 0, so all 36 layers build as type 216 and index 24+ is never read. It fires the moment tail fusion is turned on. Both the host table build and the device refresh now key off MAX_INPUTS_PER_TASK / MAX_OUTPUTS_PER_TASK. The two widths must agree -- they are the stride of the same table -- so deriving both from the TaskDesc capacity keeps them correct for any future variant. The allocation already derives from the same constants and scales automatically. Found while auditing every unbounded pointer the fused path can dereference, during the nil-address fault hunt. Not the cause of that fault. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The per-chunk attention workers write their o_acc/lse_acc partials with ordinary stores, then arrive at the chunk barrier; the last arrival runs merge_splitkv_ck_fmha over all NUM_KV_CHUNKS slots. Nothing made those partials visible to the merging worker. atom_add_release_gpu_s32 orders the arriving thread's own prior writes, but tid 0 is not the thread that wrote most of the chunk's partials -- the other 255 threads did, and __syncthreads is a block-execution barrier, not a cross-CU visibility barrier. The merging worker is a different block, so it could read a partially-written or stale slot. The failure is silent numerical corruption rather than a clean crash: merge weights each chunk by exp2(lse - m_global), so a torn lse that lands large drives every other chunk's weight to zero, and one that lands as uninitialized bits can be NaN and poison the whole attention output. Seq-len dependence comes from the empty-chunk path. At 512 tokens ntiles=32 and 14 of 30 chunks exit early without writing, so the window is narrow; at 32k ntiles=2048 gives all 30 chunks real work to write, so it is open on every chunk of every layer. Producer side: s_waitcnt + __syncthreads + threadfence_gpu before the arrival, so all 256 threads' stores are flushed. Consumer side: acquire fence + buffer_inv before merge, so it reads what they wrote rather than this CU's stale vL1 lines. Also fix the worker-state dump, which could not localize a hang: ws[3] was set to 11 once at gang entry and never updated, so with all 36 layers running inside one task every worker reported phase=11 for the whole run regardless of which layer it was stuck in. Publish 40000+ml per layer and decode it in the host dump. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
arrival counter and the write-through release flags The 32k runs wedged intermittently with every worker parked at a barrier short by exactly one arrival. Barrier-watch instrumentation caught the state: 237 workers waiting at the cross-XCD attention barrier and the QKV epoch barrier, all at the same epoch, held up by 3 stragglers a full layer behind, parked at the MoE W13->W2 barrier for two different experts. The aux capture named the mechanism. For each stuck expert the arrival counter was an exact multiple of W13_TILES -- every producer had arrived and the release had therefore fired -- yet the eight per-XCD release slots of that same expert held *different* epoch values. One producer writes all eight in one loop, so a spread across them is impossible if the writes survived. They did not survive. The release fan-out uses st_wt (sc0 sc1), which bypasses L2 and lands in HBM, while the arrival counter is an ordinary L2-resident atomic read-modify-write. Both sat in the same 64-byte line: HIER_STRIDE was 16 int32 per *expert*, with slots [0..7] the releases and [8] the counter. An L2 writeback of that line, still holding the old release values, lands on top of the fresh write-through data and reverts releases that already happened. The W2 workers for that expert then wait forever on a release that did fire. The fix gives every slot its own cache line: 10 lines of 16 int32 per expert, releases at [x*16] and the counter at [8*16]. This is the layout the sibling barrier in gang_linear_mxfp4_res_bias_rmsnorm_topk_mi300.cuh has always used -- only the MoE barrier packed everything into one line. The moe_fused_barrier tensor grows from 16*E to 160*E int32 (8KB for E=128), and MOE_BAR_* replaces the local HIER_STRIDE constants so producer and consumer cannot drift apart again. Also in this commit, the instrumentation that found it: - Barrier watch (MPK_WS_WAIT_BEGIN/TICK): records observed vs expected per spinning worker, which separates "the producer never arrived" from "the waiter expects an epoch that will never be reached". - MPK_WS_MARK: straight-line progress marks for code that is not a spin loop. Packed into a single store -- the four-store version cost 2.3x (6.99 vs 3.01 ms/token) because the buffer is pinned host memory and every store is a PCIe write, far past the point where the instrument perturbs the race it is meant to catch. - MPK_WS_WAIT_AUX: barrier-specific values, used here for the arrival counter and the agreement across the eight release slots. - The host dump now gates every watch read on the worker's live phase. The watch slots are last-write-wins and are never cleared on exit, so the first capture printed 56 "blockers" that were all idle at the scheduler while the one worker actually inside the layer went unprinted. Smoke: 2048 tokens, correct output, 3.428 ms/token -- unchanged against the 3.415 ms instrumented baseline, so the wider barrier costs nothing. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Root cause of the intermittent long-sequence hang.
The FP8 quant of the SwiGLU output issues four NT loads from one asm
volatile block with plain "=v" outputs. Because the block contains four
separate instructions, the register allocator is free to place an output
on top of an input it considers dead after the block -- and it does:
v_lshl_add_u64 v[6:7], v[4:5], 0, 16
v_lshl_add_u64 v[8:9], v[4:5], 0, 32
v_lshl_add_u64 v[10:11],v[4:5], 0, 48
global_load_dwordx4 v[4:7], v[4:5], off sc0 sc1 nt
global_load_dwordx4 v[8:11], v[6:7], off sc0 sc1 nt
global_load_dwordx4 v[12:15], v[8:9], off sc0 sc1 nt
global_load_dwordx4 v[16:19], v[10:11],off sc0 sc1 nt
Load 1's destination v[4:7] overwrites load 2's address v[6:7], and
load 2's destination v[8:11] overwrites the addresses of loads 3 and 4.
Loads 2-4 therefore address memory using whatever activation data the
earlier loads returned. Marking the outputs early-clobber ("=&v") makes
the allocator keep them disjoint; the addresses then live in v[20:27].
This single defect explains both observed failure modes. A garbage
address that faults gives the nil-address crash; one that never
completes leaves the wave parked on the following s_waitcnt vmcnt(0)
forever, which deadlocks the __syncthreads at the end of the quant and
through it the entire block -- and through the block, every worker
waiting on the layer barriers.
The capture that pinned it: with per-wave masks recorded on both sides
of the barrier, the stalled worker held wave exit mask 0xf (all four
waves cleared the W13->W2 barrier) against quant sync mask 0xe (wave 0
never reached the quant's __syncthreads), frozen identically across all
345 dumps.
Same defect fixed in the four other multi-load asm blocks that had it
(a second copy of the 4x NT load, two W2 epilogue prefetches, one OProj
prefetch). A scan of every asm block in persistent_kernel with two or
more loads and non-early-clobber vgpr outputs now comes back empty.
Also fixed here, found while reading the same function: the amax
cross-lane reduction assumed NSUBBLOCKS is a multiple of 4. It is 90 for
the W2 path (INTERMEDIATE_SIZE 2880 / 32), so the tail super-block
shuffled from two lanes whose loop condition had failed and whose amax
register was never written, giving a block scale derived from garbage.
The partner index is now clamped to the last sub-block that ran.
Instrumentation: per-wave exit/arrival masks for thread-divergent polls,
which is what made the block-split visible at all -- every previous
tracer was tid-0 only, and tid 0 was on the cleared side.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Two register hazards in software-pipelined v_mfma_scale_f32_16x16x128_f8f6f4 loops corrupt results silently -- no hang, no fault, no NaN. Each perturbs roughly one iteration out of a 22-iteration accumulation, giving 0.4-6% GEMM error that sits near the FP4 quantization noise floor. Both are present in production while the model emits fluent text at full speed. Hazard 1 (intermittent, ~17-22% of launches): the loop issues iteration N+1's ds_read into the same VGPRs iteration N's MFMA reads as sources, relying on the read landing late. That is a WAR race -- lgkmcnt tracks when LDS data reaches the VGPR, not when the MFMA finished sampling its operands, and the ISA exposes no counter for the latter. Fix is two disjoint operand banks, ping-ponged. Hazard 2 (deterministic): "s_nop 7; s_nop 0" is a 9-clock gap before v_accvgpr_read_b32, but the scaled MFMA is a 32-cycle op on CDNA4. The 9-clock tail is correct for a 4-pass MFMA and appears to be a leftover from smaller 16x16x16 f16 code. Fix is s_nop 15 x2. The test builds reference (no overlap), fixed (banked + 32clk) and broken (historical) schedules, and requires the first two to be bit-exact. The broken variant is retained so the test can prove it still detects the hazard, and warns if it ever stops reproducing -- that means the test went blind, not that the hardware got safer. Note for anyone porting the fix: with an even iteration count (22 in production) the ping-pong loop exits on a bank-1 iteration, so the trailing MFMA must read bank 0. A single-tail port silently reintroduces the same class of error. Both exit paths are spelled out in mfma_fixed. Detection only; the production loops in gang_rmsnorm_linear_mxfp4_bias_mi300.cuh and gang_moe_fused_mxfp4_mi300.cuh are not yet repaired. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
PIPELINED_QKV (gang_rmsnorm_linear_mxfp4_bias_mi300.cuh) and PIPELINED_W13_T0 / W13_T1 / W2_T0 (gang_moe_fused_mxfp4_mi300.cuh) each carried two register hazards that corrupt results with no hang, no fault and no NaN. Hazard 1 -- WAR race on MFMA sources. The loop issued iteration N+1's ds_read into the same VGPRs iteration N's MFMA reads, relying on the read landing "late". lgkmcnt tracks when LDS data reaches the VGPR, not when the MFMA finished sampling its sources; the ISA exposes no counter for the latter, and a 16x16x128 MFMA streams operands over the op's duration rather than latching at issue. When LDS returns fast the op sees mixed-iteration operands. Intermittent, ~17-22% of launches. The tell was that the B-scale alone was double-buffered through a staging register while the other four operands were overwritten in place. Fix: two disjoint operand banks, ping-pong. While the MFMA consumes bank X the prefetch writes bank 1-X, so no register is ever both a live MFMA source and an in-flight LDS destination. Bank 1 (v18, v19, v26-v29, v32-v39) added to the clobber lists. Hazard 2 -- AccVGPR read before retirement. "s_nop 7; s_nop 0" is 9 clocks (s_nop N waits N+1) before v_accvgpr_read_b32 of a 32-cycle op. Deterministic, wrong every launch. That tail is the documented wait for a 4-pass MFMA and looks carried over from 16x16x16 f16 code. Fix: "s_nop 15" x2 == 32 clocks. Measured threshold is 11 clocks (10 fails 300/300, 11 passes 300/300), corroborated by LLVM's hazard recognizer, which emits exactly s_nop 7; s_nop 2 for the intrinsic-based MFMAs in gang_linear_mxfp4_res_bias_rmsnorm_topk_kernel. Inline asm bypasses that recognizer, which is why the hand-written tail sat two clocks short. We pad to 32 anyway: the threshold is undocumented and the tail runs once per GEMM. Ping-pong makes exit parity load-bearing, so both tails are emitted. All four loops run 23 iterations (REDUCTION_SIZE 2944 / K_PER_MFMA 128), which exits via the bank-1 half with final operands in bank 0; a single-tail port would silently reintroduce the bug at the other parity. Why this survived: each hazard perturbs one iteration out of 23, bounding GEMM error at 0.4-6% -- at the FP4 quantization noise floor, so greedy argmax over ~200k logits usually picks the same token. Every check applied for weeks was a liveness check. A kernel that runs and emits fluent text is not evidence of numerical correctness. Verified in emitted ISA for all four loops (disjoint banks, both tails present) and against the regression test at both parities: ITERS=23 reference 0/19200 fixed 0/19200 broken 19200/19200 (72.04%) ITERS=22 reference 0/19200 fixed 0/19200 broken 19200/19200 (19.34%) Demo runs clean: RC=0, 3.428 ms/token at 2048 seq len, coherent output. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Until now the only correctness signal for the megakernel was greedy-token agreement against a Torch reference -- a thresholded, single-prompt proxy. It says the two paths agree or they don't, with nothing in between, so a numerical regression that bends the distribution without flipping the argmax is invisible to it. This adds a continuous, corpus-wide number. Two things were missing. The first was logits: the LM head keeps the logit row in registers and reduces it to a (max_val, abs_idx) pair, deliberately, because that removes a 393 KB write plus read every step. The second was teacher forcing -- and that one turned out to already exist. prepare_next_batch (persistent_kernel.cuh:692) only copies a sampled token into tokens[] when step + j + 1 >= prompt_len, so inside the prompt the kernel consumes whatever the host placed there. Prefill IS teacher forcing. Loading a corpus as one long prompt and running prefill-only scores every position against the reference prefix with no per-step host round-trip and no change to the megakernel loop. So only the logits sink is new code: gang_rmsnorm_linear_mxfp4_bias_argmax_mi300.cuh gains an optional logits_out_ptr plus the step to write. The store reuses the abs_idx and bias-added val the argmax path already computed, and is guarded on the pointer being non-null, so the serving path is untouched -- measured 2.437 ms/iter with PPL_MODE unset against a 2.424-2.431 baseline. The sink is attached with input_map (-1,-1,-1), NOT the (1,-1,-1) the argmax partials use. A non-negative map partitions that dim across grid_dim.x and runtime.cc pre-shifts each block's base pointer by bid.x * vocab/8; since the kernel indexes by the absolute vocab column, a pre-shifted base double-counts the offset. Getting this wrong wrote alternating 25152-column stripes from blocks 0-3 and ran blocks 4-7 off the end of the row, which scored 64.5 against a Torch 8.15 while looking like a plausible accuracy result. Both places that record the input/output split -- task_register.cc and graph.cc's task_config -- now derive num_outputs from bgraph.operators.size() instead of hardcoding 2, so runtime.cc:207's assert holds for both shapes. PPL_MODE=1 in demo.py wires it up: load a WikiText-2 slice, write the ids into tokens[], set max_seq_length so the prefill-completion stop fires exactly at the end of the slice, then cross-entropy on the host against the corpus targets. Logits land in f32, not bf16 -- bf16's ~0.4% relative precision is the same order as the GEMM error being measured. The softmax is sliced to config.vocab_size because the buffer is padded 201088 -> 201216 and the pad columns must stay out of the denominator. PPL_MODE rejects FUSE_TAIL=1 (the fused tail returns before it would reach the sink) and any --max-num-batched-tokens but 1 (the LM head GEMM consumes one row per iteration). The demo also self-checks the sink rather than trusting it: it reconstructs each of the 240 workers' owned column sets from the sink and compares against argmax_part_value/argmax_part_index. 0 index mismatches, 0 value mismatches, and sink_argmax == the token the kernel emitted -- so the sink is a faithful copy of what the in-register reduce saw, and any MPK-vs-Torch gap is accuracy, not instrumentation. 512-token WikiText-2 slice: MPK 91.84 vs Torch 36.04, ratio 2.55, with the two paths picking the same argmax at 54.2% of positions. The test gates on a ceiling and a ratio, not an exact value, because MXFP4 decode is not bit-deterministic; both gates were verified to fail when tightened. One trap worth naming, since it cost real debugging time: "top-1" is two different numbers here. Agreement with Torch's argmax (54.2%) and accuracy against the corpus targets (MPK 28.2%, Torch 39.5%) are both natural things to call top-1, and comparing one against the other across runs reads as a contradiction that isn't there. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The sink-coverage diagnostics allocated whole-tensor temporaries: `== 0.0` on [n, 201088] is a 6 GB bool at 32k and `.float()` on it is another 25 GB, on top of the 25 GB sink. That is an OOM in the instrumentation rather than in the thing being measured, so both now run in the same CH-sized chunks the scoring loop uses. run_gpt_oss_ppl_sweep.sh scores a WikiText-2 prefix at each length and summarize_ppl_sweep.py tabulates the dumps. The sweep does not use -e: the Torch reference OOMs partway up and a single length failing must not take the rest of the run with it. The summarizer reports a fixed-prefix column, and that turned out to be the whole point. Full-slice perplexity climbs steeply with length -- but it climbs on the *Torch reference* too (36.0 -> 139.1 from 512 to 1024), which no kernel change can explain. It is the corpus: WikiText-2 past position ~512 is a dense run of proper nouns averaging 6.27 nats against 3.58 for the first 511 tokens. Scoring every run over the same 511 positions removes it. MPK then reads 93.3 / 101.5 / 91.6 / 107.3 / 88.5 / 96.8 / 93.0 across 512 -> 32768 -- flat, and inside the 92-100 spread three runs at 512 alone produced. Perplexity does not degrade with context out to 32k. Latency 2.46 -> 3.70 ms/iter for 64x the context. KV chunks saturate at 30 (workers per XCD) from 4096 on, so past that each chunk just covers more KV. Both sink self-checks hold at every length including 32k: 0 index and 0 value mismatches against all 240 per-worker argmax pairs. Zero columns are isolated singletons (763 of 32767x201088 = 1.2e-7, first at row 88 column 97897) -- logits that genuinely round to 0.0f, not columns the kernel skipped. The Torch reference only reaches 1024. Its attention materializes a [64, n, n] f32 score matrix (modeling_gpt_oss.py:172) and the bf16 model already holds ~249 of the 252 GB visible, so 2048 OOMs. Long rows are MPK-only by necessity. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The sweep compared MPK against a reference running weights the megakernel does not run. MPK quantizes the LM head, QKV and O-proj -- all bf16 in the checkpoint -- down to MXFP4; the reference kept all three in bf16. (The MoE experts are natively MXFP4 in both, so they were never part of the gap.) The 2.5x perplexity ratio I reported therefore charged weight quantization to the kernel. PPL_MXFP4_MATCH=1 round-trips all three weight classes in the reference through the same quantizer, generalizing the existing head-only PPL_MXFP4_HEAD. The round-trip is verified to mutate weights ~9% and layer outputs ~10%, and to be shape-preserving at every shape used here, so the matched run is a real forward pass and not a no-op. Note USE_FP16_ACT (demo.py:1535) only covers QKV, not O-proj or the head, so it is not a substitute. USE_FP8_ACT, exported by all three CI scripts, is read nowhere in the tree at all -- dead, and worth deleting separately. On 1023 tokens the matched reference scores 100.06 against bf16's 139.08. That direction is not physical -- quantization cannot add information -- and Torch is bit-deterministic here (139.0814 twice), so it is not run noise but a small-sample artifact: MXFP4 error landing favorably on one WikiText-2 article. It means the absolute matched number should not be quoted as "the" reference until it is averaged over a much larger corpus. What does survive is the per-position comparison, which does not depend on which slice is easier. MPK is worse than the matched reference at the median position (+0.86 nats) by more than it is worse than bf16 (+0.57), and its predictive entropy is higher than both (6.27 vs 5.52 vs 4.89). So matching weight precision does NOT close the gap -- MPK is flatter and worse than a reference handed exactly the same coarse weights, and the remaining difference is kernel arithmetic, not quantization. Also corrects a claim in the previous commit: the accuracy decline across the sweep (27% -> 11%) was the corpus, not context length. Over the same 511 positions accuracy is flat at 26-28% from 512 to 32768, matching the flat fixed-prefix perplexity. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
--max-layers set a local num_layers that bounded only the MPK task graph.
GptOssModel.forward iterates self.layers unconditionally, so the Torch path
kept running all 36 layers and every MPK-vs-Torch comparison under
--max-layers was comparing an N-layer kernel against a 36-layer reference.
That is not a small skew: it reported logit correlation of 0.076 at one layer,
which reads as "the kernel is completely wrong on a single layer". With the
reference actually truncated the same number reads 0.991.
Truncating model.model.layers alongside num_layers makes the flag mean what it
says on both paths.
PPL_DUMP_LOGITS/PPL_DUMP_ROWS dump raw logit rows from either path. Derived
metrics can only say two distributions differ; the raw vectors say how, and a
scale error, a constant offset and unstructured noise are three different bugs
that look identical after a softmax.
Used together these localize the MPK-vs-Torch gap, and the answer is that
there is no defect to fix:
- per-layer divergence from a precision-matched reference is 0.077 and error
grows as sqrt(depth) (0.0766 / 0.0789 / 0.0736 at 1 / 4 / 12 layers), the
signature of independent rounding accumulating, not a systematic fault
- the logit error is zero-mean (bias/noise < 0.2 at most positions), so no
scale, index or masking error
- best-fit scale between MPK and Torch logit vectors is 1.0
- one MXFP4 GEMM costs 11.5% output error on its own -- more than MPK's
whole-layer divergence, so the kernel sits below the noise floor of the
format it runs in
- the arithmetic is MXFP4 weights x FP8 E4M3 activations on
v_mfma_scale_f32_16x16x128_f8f6f4 with f32 accumulate. Measured on a
K=3072 GEMM with this tree's own quantizers, FP8 activations contribute
0.026 relative error against MXFP4 weights' 0.210 -- 8x smaller, and the
two together give 0.212, so the activation format is not what separates
the paths
- MoE routing picks identical experts in identical order, and post-attention
RMSNorm matches to max abs diff 0.000000, so the discrete decisions agree
- no discontinuity at the sliding-window boundary (row 127 vs 140), and
correlation plateaus at 0.97-0.99 out to 512 rather than compounding
So the earlier "MPK is worse than a matched reference" claim was measuring
sqrt-depth rounding divergence between two equally-valid orderings, not a
correctness bug. Both paths land the same distance from an unquantized model;
they simply round differently, and at 36 layers that separates the argmax about
half the time on high-entropy text.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The GEMM arithmetic is not selectable. Every MI300 kernel quantizes its activations with _gang_wave_parallel_fp8_quant (FP8 E4M3, E8M0 per-128-element block scale) and feeds v_mfma_scale_f32_16x16x128_f8f6f4 against MXFP4 weights, accumulating in f32. The alternatives were removed from the kernels, but three pieces of their scaffolding outlived them. USE_FP16_ACT was the live one, and it was a bug. demo.py still honored it, packing raw bf16 into the workgroup buffer via pack_bf16_workgroup and handing it to the kernel through the same mxfp4_weight= parameter as the quantized tensor. The only guards that would have made a kernel decode those bytes as bf16 -- USE_BF16_NATIVE_WEIGHTS and USE_BF16_ACTIVATIONS -- are not defined anywhere: not in persistent_kernel.py's JIT flag list, not in CMake, nowhere. So USE_FP16_ACT=1 fed bf16 to the MXFP4 unpacker. The bf16 buffer is larger than the MXFP4 one it replaced (368640 vs 97920 bytes per QKV workgroup), so there was no OOB read and no fault: it decoded bf16 bit patterns as E2M1 nibbles and produced confident garbage. Silent wrong answers, no crash. USE_FP8_ACT was harmless but misleading: exported by all three CI scripts and read by nothing, and documented in the demo README as a togglable "benchmark config" for behavior that is now unconditional. USE_FP4_ACTIVATIONS in gang_full_layer_fused_mi300.cuh selected a half-width O-proj LDS staging area. Never defined, so the FP8 sizing always won; inlined with a comment naming the format instead. Verified: demo runs clean at 2.463 ms/iter (baseline 2.42-2.47), and the 512- token perplexity dump is 97.7 with the sink self-check passing on all 240 workers -- inside the known run-to-run MXFP4 spread (81.2% argmax agreement against the committed baseline matches the 81-83% cross-run determinism of this kernel). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
… arrival
f1fa720 ("widen QKV epoch barrier to all workers") fixed a real intermittent
deadlock but cost ~0.27ms/iter (2.19 -> 2.46). Bisected by header-only revert
onto the current tree; every other candidate (278bb7d, 4b393d8, d14b40b)
measured inside noise.
The cost is not the wait. Device timers (MPK_DEVICE_TIMING=1) put qkv_bar at
~2.5us -- blocking at Phase 2 is cheap. The actual mechanism is a lost
prefetch overlap: before f1fa720, ranks >= NUM_KV_CHUNKS skipped Phase 2
entirely and ran ahead to Phase 6, where the O-proj weight buffer_load_lds is
issued *before* the attn_release poll precisely so the DMA overlaps that
spin-wait. Widening the barrier pinned those workers at Phase 2 and serialized
their DMA behind QKV, x36 layers.
Split arrival from waiting. All workers_per_xcd still arrive -- arrival is what
carries the ordering the deadlock fix depends on, and both properties hold on
arrival alone:
1. the bump is still gated on all arrivals, and every poller still snapshots
before arriving, so no bump can land between a poller's snapshot and its
arrival;
2. a non-polling worker still cannot reach Phase 6/7b before arriving, and
the epoch cannot bump until it does, so its attn_release / routing_ready
"+1" snapshots still happen-before their producers.
Only ranks < NUM_KV_CHUNKS -- the workers that actually read this layer's QKV
output -- block on the epoch.
Also fixes a pre-existing build break: the O-proj/TopK call passed a _ts_base
that was declared nowhere, so MPK_DEVICE_TIMING=1 did not compile. Nothing
reads those slots (the [FUSED_PHASE] printf derives everything from
_fused_t0.._fused_t4), and the sibling task omits the argument, so it is now
nullptr. This is what unblocked the attribution above.
Measured, GPT-OSS 120B bs=1 seq 512:
before 2.461 / after 2.383, 2.386, 2.394, 2.400 ms/iter
seq 4096 stress: 6 runs x 4094 iters, no hang or fault
perplexity 92.1 / 95.7 / 99.4 vs baseline 94.3 / 95.1 / 112.5 -- inside the
baseline MXFP4 run-to-run spread, and tighter
sink/argmax self-check: 0 index, 0 value mismatches over 240 workers
run_ci_tests_gpt_oss.sh: pass (longest_common_block=15, need >= 8)
Recovers ~1/4 of the gap; the remaining ~0.24ms is not yet explained.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The fused-layer task computed its three barrier release values by snapshotting "current counter + 1" at task entry. That is only valid if the reader is ordered before this layer's producer, and nothing guaranteed it: all 36 layers run in one task (the ml loop in persistent_kernel.cuh) with only a per-block __syncthreads between them, so workers skew freely across layer boundaries. A worker a full layer behind could read an already-bumped counter and then wait for a bump this layer never produces. f1fa720 worked around this by making every worker on the XCD arrive at the Phase 2 barrier, which supplied the missing ordering. It also cost 2.19 -> 2.46 ms/iter; 2c1071c recovered most of that by splitting arrival from waiting. Publish (pc_iter - 1) * ml_num_layers + ml into the free int32 of the n_tile union member and derive all three expected values from it. All three counters start at 0, bump once per layer, and are never reset, so the value each reaches is a pure function of the layer index -- every worker computes it identically with no shared read and no arrival requirement. The Phase 2 participant set narrows back to the workers that actually need the barrier. This is a correctness change. Its latency effect measured neutral (2.366-2.404 vs 2.386 before), which falsifies the Phase 6 attribution that motivated the narrowing -- MPK_DEVICE_TIMING=1 cannot localize the remaining ~0.2ms, because ~147k printfs inflate iterations to ~56ms and xcd_barrier's median then moves opposite to real latency. Verified: 12,282 iterations at seq 4096 across 3 runs, no hang or fault; CI longest_common_block=18 (need >= 8); perplexity 97.9/100.9/98.7, inside the 94.3-112.5 baseline spread. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The MPK_WS_* phase/barrier breadcrumbs were always-on, guarded only by a runtime `g_ws_dev != nullptr` check. g_ws_dev is pinned *host* memory, so each store is a PCIe write, and there are ~30 sites across the two hottest task headers -- several inside spin loops, where the store sits on the poll path. Put them behind a compile-time flag, matching the MPK_NIL_TRIPWIRE convention, and add the MPK_WORKER_STATE env var so the dump stays reachable: it is how the fused-layer deadlocks were attributed. 2.386 -> 2.321 ms/iter at seq 512. The 11 inline `precomp_dbg_worker_state != nullptr` sites in the dispatch loop were gated the same way but measured neutral -- the compiler had already hoisted that check. This also explains part of a measurement I had been chasing. The 2.19 figure for "reverted f1fa720" came from restoring the whole fused header to f1fa720~1, which predates beac230 and so stripped these breadcrumbs too. Some of what I attributed to the Phase 2 barrier was always this. Verified: CI longest_common_block=18 (need >= 8); MPK_WORKER_STATE=1 still compiles and runs (9.79 ms/iter, showing what the PCIe stores cost when the buffer is live). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
db48239 fixed a real race -- chunk workers write o_acc/lse_acc partials
with ordinary stores, and nothing made them visible to the merging
worker -- but fixed it at agent scope. The barrier it guards is per-XCD
(chunk_barrier[xcd_id * 16]), so producers and consumer always share one
32MB L2, and agent scope pays for cross-XCD coherence that no consumer
here can observe.
On gfx950 (ISA verified):
release agent -> buffer_wbl2 sc1 ; s_waitcnt vmcnt(0)
acquire agent -> buffer_inv sc1
either at workgroup scope -> no instruction at all
So the old sequence flushed the whole L2 to HBM on every chunk worker on
every layer, then had the merger invalidate that same L2 -- discarding
lines this XCD's own chunk workers had just written and forcing a re-read
from HBM. mpk_atoms.cuh already documents why that is unnecessary: "all
CUs within an XCD share the same 32MB L2 ... No buffer_wbl2 required".
Keep exactly the part that carries the guarantee:
producer: s_waitcnt vmcnt(0), retiring all 256 threads' stores into L2
before tid 0's release atomic
consumer: buffer_inv (no sc1), dropping stale per-CU vL1 lines
Workgroup scope cannot substitute -- it emits nothing, and the consumer
is a different block.
2.321 -> 2.175 ms/iter at seq 512. For reference, deleting the fences
outright (incorrect) measured 2.161, so nearly all the cost was the L2
traffic rather than the ordering.
Verified in the regime where this race is actually reachable. db48239
notes the window is narrow at 512 (14 of 30 chunks exit early) and open
on every chunk of every layer at 32k:
seq 32768: 32,766 iters, no NaN/inf, no fault
seq 8192: 8,190 iters, coherent output, no NaN/inf
CI: longest_common_block=18 (need >= 8)
Also drops the explicit s_waitcnt vmcnt(0) that db48239 placed before the
__syncthreads: the fence after it already drained, so it was a second
redundant stall of all 256 threads ahead of the barrier.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
c9d1e94 gated the MPK_WS_* call sites but left the pointer publication at kernel entry unconditional -- a store plus a __syncthreads per worker per launch, for a pointer nothing reads when MPK_WORKER_STATE is off. Wrap it in the same #ifdef the tripwire publication directly above already uses. 2.175 -> 2.155 ms/iter at seq 512. Swept the rest of the device path while here; everything else is already gated and needs no change: - MPK_TW_SUB / config.tripwire sites -> MPK_NIL_TRIPWIRE - MOE_DBG_ENTRY / MOE_DBG_SUBPHASE -> MPK_NIL_TRIPWIRE - g_subphase_scratch timestamps -> MPK_ENABLE_MOE_SUBPHASE, or #if 0 - [FUSED_PHASE] printfs -> MPK_ENABLE_DEVICE_TASK_TIMING - the MPK_WS_ON(config) dispatch sites -> compile to (false), body deleted - the host-side hipHostMalloc -> already env-gated, nullptr default Verified: CI longest_common_block=18 (need >= 8); MPK_WORKER_STATE=1 still prints "worker-state tracing ON (240 workers)" and runs. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Three prints fired on every default run and nothing parsed them: [FWD_PASS_CURVE] 10 lines/run [MPK] XCD templates uploaded 1 line/run [MPK] ML table validation 1 line/run The decile curve was a scaling diagnostic that [FWD_PASS_TOTAL] and the per-iteration [FWD_PASS] lines already cover between them; it read the same decimated ring, so it added no information the log did not have. The template-upload line reported a value the caller already knows. The ML table validation loop stays -- it names the exact slot behind a nil-address GPU fault, which has no usable backtrace otherwise, and the per-slot ML_NULL_IN/ML_NULL_OUT prints only fire when a slot is actually null. Only the summary is now conditional on null_in || null_out: a healthy graph printed "0 null inputs, 0 null outputs" every run, which trains the reader to skip the line that matters. These all fire once at terminate or at startup, not per iteration, so this is a log-cleanliness change and not a latency one. Measured to confirm: 2.164 ms/iter at seq 512, unchanged from 2.155. g_fwdpass_stride/_count, max_tpw and total_template_entries all remain live after the removals -- no orphaned state, no unused-variable warnings. Also .gitignore the artifacts these runs leave behind: outputs/, the tests/standalone binaries built by build.sh, and hipcc -save-temps intermediates. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The Clang format CI job checks whole files, so any file we edited had to come out clean under the repo's .clang-format at version 15 (matching .github/workflows/code-format.yml). Ran clang-format 15.0.7 locally -- same version the action pins -- and applied its own output verbatim. The change is larger than the lines we wrote. clang-format's comment alignment and argument packing are context-sensitive, so inserting a block can pull adjacent untouched lines into a different alignment group; 37 of the reformatted lines exist verbatim upstream and were clean there. Nothing to do about that short of not formatting the file. Verified no semantics changed: comparing comment-stripped, whitespace- stripped token streams before and after, every file is token-identical except one declaration-specifier reorder that clang-format performed itself -- -static volatile bool g_tw_snap_stop = false; +static bool volatile g_tw_snap_stop = false; which is equivalent in C++ (decl-specifiers are order-independent). tests/standalone/test_mfma_pipeline_hazards.hip is ours and is reformatted here too. Because it is a regression test whose inline-asm constraint lists got re-wrapped, it was rebuilt and rerun rather than assumed: 19200/19200 blocks still differ for the known-broken schedule (worst rel err 73.39%), fixed and reference still bit-exact. It still detects the hazard it exists to detect. The four pre-existing tests/standalone/test_ck_*.hip and test_mfma_simple.hip files are left alone. They violate the same check at upstream HEAD (61-90 warnings each) and predate this branch; reformatting them here would bury this diff in unrelated churn. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Correctness, performance, and measurement work on the MPK GPT-OSS-120B path for
MI350/MI355 (gfx950). 24 commits on top of
c33065f.Six deadlocks and two silent numeric-corruption bugs fixed, at no latency
cost. The deadlock fixes themselves cost 2.19 → 2.46 ms/iter; the perf commits
here recover that, ending at 2.16. Net this is roughly latency-neutral — the
value is the correctness, not a speedup.
Correctness
Six hangs, each reproduced and root-caused:
so workers waited on chunks that were never dispatched.
against their producers.
line with the write-through release flags.
"=v"→"=&v"). Oneroot cause behind both a deadlock and a nil-address GPU fault: the compiler
reused an output register as an input address across the asm block.
Two silent gfx950 scaled-MFMA pipeline hazards: no hang, no fault, no NaN —
roughly 1 iteration in 23 had GEMM results off by 0.4–6%. Liveness checks cannot
catch this class; found by diffing against a serialized reference. Covered by
tests/standalone/test_mfma_pipeline_hazards.hip.Also: fences where routing and attention data crossed a barrier without
ordering, and host-side validation of the multi-layer pointer tables, which
names the exact null slot behind a nil-address fault (that fault otherwise has
no usable backtrace).
Performance
Two causes, bisected one commit at a time — neither is the widened barrier, so
part of what looked like a barrier regression was instrumentation that landed
alongside it:
MPK_WS_*breadcrumbs had no#ifdef, only a runtime null check, at ~30 sites in the two hottest headers— several inside spin loops, storing to pinned host memory. Now behind
MPK_WORKER_STATE.agentscope, but the barrier itguards is per-XCD, and on MI300/MI350 all CUs in an XCD share one L2. Agent
scope emitted
buffer_wbl2 sc1(a full L2→HBM flush) per chunk worker perlayer, and the merger then
buffer_inv sc1'd that same L2 — discarding linesits own XCD had just written. Scoping to the XCD keeps the ordering and drops
the flush. Fence lowering was verified against real gfx950 ISA.
Barrier arrival carries the happens-before ordering these counters depend on;
the wait does not. Narrowing who blocks is safe, narrowing who arrives is not
— unless the expected values stop being snapshots, which the deterministic
layer index makes true.
Perplexity
The LM head reduces logits to an argmax in registers and never writes them to
HBM, so the MPK path had no way to measure output quality. Adds an opt-in
PPL_MODElogits sink (off by default, so the serving path is unchanged), a512–32768 sweep harness, and a CI test. Prefill is already teacher forcing —
the device writeback is guarded by
step + j + 1 >= prompt_len— so aprefill-only run scores a document with no change to the megakernel loop.
Perplexity at seq 512 has a wide run-to-run spread under MXFP4 (94.3–112.5
observed), so single-run differences are not meaningful.
Verification
longest_common_block=18against the Torch reference (threshold 8) on everycommit in the series.
matches before, 4/4 after.