fix(train): gfx1201 ROCm fixes + unsloth bnb-4bit Gemma 4 31B QLoRA

[ulises-c]

Closes #100

Summary

• Patch caching_allocator_warmup on HIP/ROCm to suppress "Found no NVIDIA driver" raised by the CUDA-only allocator warmup in recent transformers
• Force device_map={"": 0} for the whole QLoRA path so bitsandbytes can place all layers on the single GPU ("auto" can dispatch layers to CPU, which bnb rejects). Covers both live quantization and the pre-quantized dynamic 4-bit unsloth checkpoint (which keeps embeddings/lm_head/gates in 16-bit)
• Target the community pre-quantized unsloth/gemma-4-31B-it-unsloth-bnb-4bit checkpoint (issue Investigate gfx1201 (RDNA4) training: Flash Linear Attention + Gemma4-31B QLoRA OOM #100): downloads already-NF4 (~19 GB), so it skips the ~62 GB BF16 CPU staging at load that was the R9700's actual blocker — no L40S prequant + rsync needed
• Add make train-gemma4-31b-stage2-unsloth (TRAIN_PREQ=true + unsloth base) and document the prequant vs live-quant launch paths in train-sft-stage2-gemma4-31b.env
• Drop the 26B-A4B config (wrong target; 31B is the real one)

Test plan

• --dry-run validates the unsloth 31B config plumbing (config parse, LoRA regex, SFTConfig, data load) on the dev box
• make test-gpu-stack passes on gfx1201 (confirms bnb 4-bit runtime) — already 13/13 per Investigate gfx1201 (RDNA4) training: Flash Linear Attention + Gemma4-31B QLoRA OOM #100
• make train-gemma4-31b-stage2-unsloth loads without "Found no NVIDIA driver" on ROCm
• Loads without bitsandbytes CPU-dispatch rejection (the device_map fix)
• Gemma 4 31B QLoRA trains end-to-end on R9700 from the unsloth checkpoint

Note: dry-run only validates plumbing. Model load, device_map, and the caching_allocator_warmup patch require the AMD gfx1201 box (this dev machine is NVIDIA).

HuggingFace Model

Gemma4-31B-it-bnb-4bit

🤖 Generated with Claude Code

[codecov]

Codecov Report

❌ Patch coverage is 87.87879% with 24 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
src/project/hf_callback.py	88.00%	9 Missing and 6 partials ⚠️
src/project/dataset.py	86.76%	7 Missing and 2 partials ⚠️

📢 Thoughts on this report? Let us know!

[ulises-c]

gfx1201 learnings from PR #99 — relevant to this PR

Picking up Gemma 4 31B QLoRA on the R9700 after the Qwen3.5 classifier work in #99. Key findings that apply here:

1. hipBLASLt may be safe for Gemma 4 — worth testing first

The GPU page faults (rc=134) in #99 were specific to torch_chunk_gated_delta_rule in Qwen3.5's linear attention layers. Gemma 4 uses standard softmax attention — no linear attention, no delta rule op. The crash was not a general hipBLASLt instability, just that one kernel path. Recommend testing with TORCH_USE_HIPBLASLT=1 on a clean GPU before defaulting to the rocBLAS fallback. If it holds for 20+ steps, the 4.1× matmul speedup is available (measured: 5.2ms vs 21.5ms on gfx1201).

2. Add PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True to the launch env

This was what finally stabilized hipBLASLt past the early crash zone in #99 — it prevents VRAM fragmentation during allocation. For a 31B NF4 load (~19 GB weights + optimizer + activations), fragmentation during loading is a real risk. Add it to train-sft-stage2-gemma4-31b.env alongside the other env vars.

3. GPU state corruption cascade — verify clean before each attempt

Every GPU page fault (rc=134) on gfx1201 leaves the KFD in a dirty state. Subsequent runs — even with the correct config — also fault at random steps until you kill -9 <pid> and confirm rocm-smi --showpids returns zero processes. This caused repeated silent failures in #99. The test plan checklist should include a GPU-clean check before each run:

rocm-smi --showpids | grep python && echo "GPU dirty — kill before retrying" || echo "GPU clean"

4. Use setsid instead of nohup for the training wrapper

nohup bash script.sh & doesn't create a new process group — children survive kill -9 -- -$PGID. setsid bash script.sh & puts the entire tree in its own group. Kill the whole tree cleanly with kill -9 -- -$PGID. Prevents the overlapping-process-on-same-GPU issue that caused multiple spurious crashes in this session.

5. test_training_gpu.py from #99

scripts/test_training_gpu.py (added in #99) runs a 7-step GPU smoke test covering fp32/bf16 matmul, sdp-attn, linear-attn proxy, grad-ckpt compat, and VRAM headroom. Useful as a pre-flight before a multi-hour QLoRA run. The linear-attn proxy isn't relevant for Gemma 4, but the matmul/sdp-attn/VRAM steps are.

[ulises-c]

Notes from #94

# From 5090 Run
Critical finding - the SFT model is generatively broken.

The model generated 2048 tokens in 38s (~53 tok/s, fast) but output is textbook repetition collapse:

这样可以帮助他建立更强的数学基础。这样可以帮助他建立更强的数学基础。

这样可以帮助他建立更强的数学基础。［~70+ repetitions, never emits EOS］

Going back to the morning's smoke test, it ALSO showed early degeneration patterns - the trailing `*（此处是观察点....）* *（此处是引导点....）* *（此处是总结点....）*` was the same loop just truncated at 200 tokens. I missed it.

Diagnosis: the model never emits EOS. Without a max_tokens cap, every teacher call runs to context Limit (~150K tokens, taking ~45 min each). The 88x stowdown I saw wasn't cold-start overhead - it was the model never finishing a single response.

What this means for the eval

We sit somewhere in the STATUS_REPORT §2.8 outcome matrix's "fine-tuned Gemma loses both" cell:

> Either undertraining or data-format issue / Document honestly in §Limitations; Locked headline stands at 72.24

Two paths forward

Path A - Run the eval with the cap, report the (likely bad) numbers as evidence of the failure mode. Negative result is still a paper contribution: documents that Pattern-A SFT on this size dataset caused output collapse. Locked headline 72.24 stays as A-grade submission.

Path B - Diagnose root cause first. Most likely candidates:

1. Pattern-A format trained the model to depend on specific structural markers - without them at inference,
2. 05_K_M guantization damaged a small model - test by re-running smoke on the BF16 merged checkpoint (but won't fit on 5090 VRAM; CPU inference is very staw)
3. Chat template mismatch-training used a patched template; serving uses llama.cpp's llama.cpp's Possible structural mismatch.

Which way do you want to play it? My recommendation is Path A first (~45 min to get hard numbers showing the failure) then Path B (post-mortem after we have the data). The negative-result narrative is paper-worthy and matches what §2.8's outcome matrix already anticipates.

[ulises-c]

Session update — schema drift fix + training paused

Root cause (from 5090 run)

The 5090-trained adapter is generatively broken: never emits EOS, runs to context limit (~150K tokens, ~45 min/call). Smoke test and full eval showed the same repetition collapse. Root cause diagnosed via A/B (same model, same GGUF, same server):

• Training-shape prompt → 20 tokens, EOS, clean Socratic question (0.5 s)
• Inference-shape prompt → 2048-token cap, 这样可以帮助他建立更强的数学基础 × 70+ reps (38 s)

The failure is a train/serve schema mismatch, not quant damage or chat-template corruption:

Axis	Training (old socrat-zh/en)	Inference (socrates_teacher())
System prompt	1-line _SOCRAT_ZH_SYSTEM + problem context	7-line rules block, no problem context
Message structure	Multi-turn (N user/assistant pairs)	Single user message
History	Actual HF message turns	历史对话记录:\n学生: ...\n老师: ... blob
Current input label	None	当前学生输入: prefix

Fix: new socrat-zh-sft / socrat-en-sft sources

Added to src/project/dataset.py — produce one (system, user, assistant) triple per dialogue turn in the exact format socrates_teacher() constructs at inference:

system: [7-line rules block — identical to socratic_teaching_system.py:404-415]
user:   \n历史对话记录:\n{学生/老师 pairs}\n\n当前学生输入: {s}\n\n苏格拉底教学顾问评估结果: 学生处于 {state} 状态\n苏格拉底教学顾问建议的操作: {action}\n
asst:   {teacher turn}

Render-diff result: MATCH: True (system and user byte-identical to inference for a real dialogue turn).

Backward-compatible: socrat-zh / socrat-en multi-turn loaders are unchanged — build_dpo_pairs.py and tests are unaffected.

Other changes this session

• quantization_config=None bug: explicit None passed to AutoModelForCausalLM.from_pretrained caused auto_factory.py to overwrite the unsloth checkpoint's config.json quantization config with None → supports_quant_method crash. Fixed by only passing quantization_config when it's not None (same for dtype).
• W&B: wandb>=0.18.0 added as a core dep; _check_wandb() auto-detects login via ~/.netrc, warns and disables tracking if not authenticated. report_to=["wandb"] wired into SFTConfig. Run name defaults to output-dir basename.
• .gitignore: wandb/ local run artifacts added.

Dry-run

Sources: ['socrat-zh-sft', 'socrat-en-sft']
train: 77,202 records  (per-turn; was 12,244 per-dialogue)
eval:   8,578 records
=== Dry run complete — all checks passed ===

Status

Training paused — running consult before re-launch. Blocking gate: post-checkpoint generation test on inference-shape prompt must confirm clean EOS before trusting the full 55h run.

[ulises-c]

Code review patch — dialogue-level split fix (d04bc83)

Finding: the initial socrat-zh-sft / socrat-en-sft loaders called _split on the flat per-turn list. Turn N of a dialogue could land in train while turn N+1 landed in test; test records then contained turn N's teacher response verbatim in their 历史对话记录 blob — contaminating the eval split with training targets.

Fix: split at the dialogue level first, then expand to per-turn records. All turns of a dialogue now stay in the same partition, matching the behaviour of the original multi-turn loaders.

# before — splits flat per-turn list
converted = [turn for r in raw for turn in _socrat_zh_to_sft_records(r)]
return _split(converted, split, seed)          # ← scatters a dialogue's turns

# after — splits dialogues, then expands
split_raw = _split(raw, split, seed)           # ← whole dialogue goes to one split
return [turn for r in split_raw for turn in _socrat_zh_to_sft_records(r)]

Record counts unchanged: 77,202 train / 8,578 eval. All 143 tests pass. Pushed as d04bc83.

Other review findings (accepted / deferred):

• Evaluation string drift (学生处于 {state} 状态 vs free-form consultant text) — accepted residual; the A/B test confirmed this terse form is sufficient to elicit clean EOS. Filed for future work when a re-annotated dataset is available.
• _build_inference_user_message divergence risk — logged; mitigation is the render-diff gate before any future change to get_formatted_history().
• _check_wandb always-on — accepted; W&B is now a core dep and credentials are present. A WANDB_DISABLED=true env var can suppress it if needed.

[ulises-c]

Bilingual (EN/ZH) teacher scaffolding — pros/cons of building it out

Considered making the SFT format language-aware (English scaffolding for socrat-en-sft, Chinese for socrat-zh-sft) rather than the current Chinese-only block applied to both. Full buildout plan tracked in #108. Summary of the trade-offs:

What it concretely requires

It's bigger than the teacher prompt alone, because several pieces are Chinese-only:

1. English system block (translate the 7-line rules at socratic_teaching_system.py:404-415).
2. English canonical state→action map (all 35 entries). The EN dataset's action strings mirror the Chinese dataset annotation, not the canonical map (a1 → 生成一个与解题相关的子问题), so this is new authored content.
3. English user-message labels (历史对话记录/当前学生输入/eval/action/学生/老师).
4. Language-aware refactor of socrates_teacher / get_formatted_history / get_action_for_state + a language param threaded through serving.
5. The hidden blocker — the consultant. evaluation/state come from a Chinese-only consultant. Real English parity needs it to classify on English dialogues and emit English evaluation prose; until then English serving still yields a Chinese eval line, so teacher-side bilingual alone doesn't fully close the loop.

Pros

• A genuinely bilingual product — English students get English Socratic teaching; stronger paper story than English content under Chinese scaffolding.
• Makes the SocratDataset-EN split coherent (training English content with Chinese scaffolding is itself a mild drift).
• The shared-source-of-truth refactor permanently kills the drift bug class behind Stage 2b SFT: GGUF pipeline shipped, eval blocked by train/serve schema drift #94.

Cons / risks

• Scope creep vs. the Stage 2b SFT: GGUF pipeline shipped, eval blocked by train/serve schema drift #94 fix. Stage 2b SFT: GGUF pipeline shipped, eval blocked by train/serve schema drift #94 is "stop the collapse," already achieved by the Chinese-only fix (canonical action map + free-form evaluation field + structural match). Bilingual is a new feature, not a fix.
• The consultant lift is the expensive ~80% and must be scoped/validated separately.
• Re-baseline risk on the locked 72.24 / Tables 6/14 — any socrates_teacher edit is sensitive. zh should stay byte-rendered-identical (the Gemma template trims, so boundary whitespace vanishes), but that must be empirically re-verified before trusting it.
• Translation quality — I'd be authoring the canonical English rules + actions; needs a fluent reviewer.
• Doubles GPU validation (separate English EOS-gate + eval run) on top of the priority zh 55h run.

Recommendation: phase it

Ship the Chinese-only fix first (it's the real #94 fix), launch the zh run, and treat bilingual as a deliberate follow-up — and when doing it, do the consultant too, or English parity is only half-built. Full detail, design, and acceptance criteria in #108.

Note (infra): uv run pytest fails in this env — pytest isn't a console entry point; use uv run python -m pytest. Worth checking the post-test hook / CI doesn't assume the former.

[ulises-c]

Phase 0 applied — Stage-2b SFT is now dev-box "100% ready" (eb12dbd)

External-consultant review + patch finishing the #94 train/serve schema fix that closes out this PR. Phase 0 removes the last two drifting lines between the socrat-zh-sft / socrat-en-sft training records and the prompt socrates_teacher() actually sends at inference — the residuals that, left in, would have re-invited the original non-terminating-repetition collapse.

What was drifting, and the fix

Line	Was emitting	Inference actually sends	Fix
苏格拉底教学顾问建议的操作:	raw dataset turn["action"]	get_action_for_state(state)	loader now calls the same canonical map — byte-exact
苏格拉底教学顾问评估结果:	templated 学生处于 {state} 状态	live consultant prose	loader now uses the dataset's per-turn free-form evaluation field

The action drift was systematic — a1 (every dialogue's opening teaching turn) never matched (生成一个问题 vs 生成一个与解题相关的子问题), plus several c-states.

Changes

• socratic_teaching_system.py — hoisted state_to_action to a module-level STATE_TO_ACTION + get_action_for_state() so the loader imports the exact inference source of truth. Behavior-preserving: AST-verified all 35 entries identical to HEAD, and the live socrates_teacher render is byte-for-byte unchanged for a known turn (no eval re-baseline — headline 72.24 / Tables 6/14 untouched).
• dataset.py — _build_inference_user_message takes evaluation; both -sft builders source the action from get_action_for_state(state) and the eval from the dataset evaluation field. Target-turn gating still keys on dataset state+action presence, so the record set is unchanged.
• test_sft_inference_format.py — un-rigged the action dimension (the parity test previously passed the dataset action straight into inference, so it never tested the action line). The render-diff gate now asserts the action line render-equal too; the eval line is isolated as the sole expected residual.

Verification (this dev box)

• ✅ Full suite 154 passed / 2 skipped; ruff / pyright / codespell / shellcheck clean (pre-commit gate).
• ✅ Loaded the real HF sources — socrat-zh-sft and socrat-en-sft, 42,890 records each — and confirmed 0% template-fallback evals: every record carries the dataset's free-form evaluation, so the eval-line fix is not a silent no-op.
• ✅ Official make train-gemma4-31b-dry-run plumbing passes.
• ✅ zh inference prompt confirmed identical to pre-change HEAD.

⚠️ The EOS gate on the R9700 is still MANDATORY — and must use a live free-form eval

Phase 0 closes the structural and action-line drift, but the eval line still drifts at real serve time — the consultant regenerates fresh prose on every call, so it will not match the dataset's stored annotation token-for-token. test_evaluation_line_is_the_only_residual_drift_under_freeform_consultant pins this as the irreducible residual. Do not read "parity test passes" as "model is collapse-proof." Before trusting the multi-hour run, train a short checkpoint and prove EOS termination on a prompt built with a realistic multi-sentence Chinese evaluation, end-to-end through process_student_input (gold) or via the _capture_inference_prompt path (lighter). A gate built on the 学生处于 {state} template is blind to exactly this residual and will pass while shipping a broken model.

Carry-forward notes for the R9700 launch (§5)

• Recompute run size — the config comment ("~2298 steps / 11 GB") is stale (per-dialogue era). Per-turn is ~77k train records → ~14k steps at bs1×ga16×3ep (multi-day). With TRAIN_SAVE_STEPS=100 that's 140+ adapter checkpoints (~70 GB) — raise TRAIN_SAVE_STEPS and/or set save_total_limit before launch.
• GPU must be clean — rocm-smi --showpids empty (a prior rc=134 leaves the KFD dirty).
• hipBLASLt off by default (safe); =1 is an optional ~4× perf bet — test 20+ steps first (feat(train): Qwen3.5-0.8B LoRA classifier — W&B tracking + transformers 5.9 drift fix #99).

Out of scope / known gaps (deferred, not bugs)

• Bilingual (EN/ZH) Socratic teacher: language-aware train + serve scaffolding #108 (bilingual EN/ZH) — socrat-en-sft stays Chinese-scaffolded: its English evaluation prose is dropped into the Chinese teacher prompt as-is. Intended for now.
• Parity-test blind spot (pre-existing, not introduced here): the action line is now structurally tautological (both train + serve route through get_action_for_state, so map coverage cannot break parity — good). But that means a future edit corrupting a non-a1 map entry (e.g. b2) would corrupt both sides identically and pass every test. Only a1 is pinned by an independent literal assertion; the 35-entry AST-vs-HEAD purity check was point-in-time, not CI. Flagging as a known gap, not fixing in this PR.

[ulises-c]

gfx1201 ROCm training — issues encountered on feat/gfx1201-rdna4-qlora-fla-training

Attempting to reproduce the EOS-gate checkpoint run on the R9700 (gfx1201, 32 GB). Documenting what went wrong and what fixed it.

Issue 1 — make test-vllm auto-installed vLLM and broke the venv

The test-vllm Makefile target detected vLLM was absent and auto-installed vllm==0.22.0, which pulled in torchvision==0.26.0 (CUDA build). That torchvision registers a fake torchvision::nms op at import time that doesn't exist in the ROCm torch, breaking peft and transformers imports:

RuntimeError: operator torchvision::nms does not exist
ModuleNotFoundError: Could not import module 'BloomPreTrainedModel'

Fix: uv pip uninstall vllm torchvision. Both packages are inference-only; the training stack (peft, trl, transformers) doesn't need them.

Recommendation: The test-vllm script should check for vLLM but not auto-install — or install into an isolated env. Auto-installing into the shared training venv is destructive.

Issue 2 — expandable_segments:True unsupported on HIP, caused GPU page fault at step 16

PYTORCH_HIP_ALLOC_CONF=garbage_collection_threshold:0.8,expandable_segments:True generated a warning at startup:

UserWarning: expandable_segments not supported on this platform
(Triggered internally at /pytorch/c10/hip/HIPAllocatorConfig.h:40.)

With this set, the run faulted at step 16 every time:

Memory access fault by GPU node-1 on address 0x7f78cbc00000. Reason: Page not present or supervisor privilege.
GPU core dump failed

Fix: Stripped expandable_segments:True from all Makefile targets. The note in #101 references PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True — on ROCm the variable is PYTORCH_HIP_ALLOC_CONF, and expandable_segments is not a supported key for the HIP allocator.

Issue 3 — GPU KFD dirty state on repeated crashes (per #101 note §3)

After each page fault the KFD is left in a dirty state. Added rocm-smi --showpids check before each retry. Both subsequent retries were preceded by confirming a clean GPU state.

Current state

Running with TORCH_USE_HIPBLASLT=1 and setsid (per §1 and §4 of the #101 notes). Past step 16 is the validation point — will update once we confirm stability past that threshold.

VRAM at step 10 (last clean run): alloc=21.1 GB, reserved=27.8 GB / 31.9 GB total — within headroom, not OOM.

[ulises-c]

Issue posted in rocm-libraries #7992

[ulises-c]

Running the crawl monitor (latest) + crash diagnosis

The monitor now posts to a single self-updating comment on the Training Log issue (#120) instead of spamming this PR per crash. Progress rows land every 50 steps (time / loss / grad_norm / lr); crash, STALLED, and COMPLETE are rows in the same table. Landed in e9ae500.

1. Get the latest monitor onto the training host

The host keeps running whatever version was launched — editing/pulling does not change a running monitor. So pull, then relaunch:

cd <repo-on-host>
git pull                                  # must include e9ae500
pgrep -af 'monitor_stage2\.sh'            # is an OLD one still running? kill it first:
# pkill -f 'monitor_stage2\.sh'

2. Launch it

nohup bash scripts/monitor_stage2.sh > outputs/monitor_stage2.log 2>&1 &
tail -f outputs/monitor_stage2.log

It launches training itself (make train-gemma4-31b-stage2-unsloth) with a fresh TRAIN_DATA_SEED, polls every 5 min, and on a fault: archives the log + dmesg, cleans the GPU, quarantines a partial checkpoint, and resumes. A ▶ monitor start row in #120 confirms it's posting.

Every launch is GPU-gated — including the first. Before train_sft.py starts, the monitor (a) waits up to 180s for the GPU/KFD to be verified clean (test_gpu_stack.sh --wait-clean, so it won't fault on stale PTEs), then (b) runs gpu-preflight (test_gpu_stack.sh --preflight), which is a Makefile prereq of the train target. If preflight fails, make aborts and you'll see launch aborted (gpu-preflight failed?) — counts as no forward progress in monitor_stage2.log. You don't run make gpu-preflight yourself — it's automatic.

New crawl (not a resume)? Create a fresh placeholder comment on #120 and paste its numeric id into LOG_COMMENT_ID in monitor_stage2.sh (same "bump per run" convention as WANDB_RUN_ID).

3. If it crashes and nothing posts — diagnose on the host

Nothing posting almost always means the monitor process wasn't running when training faulted (a silent crash), not a bad fault. Run:

cd <repo-on-host>
pgrep -af 'monitor_stage2\.sh'                          # monitor alive?
pgrep -af 'train_sft\.py'                               # training alive?
tail -40 outputs/monitor_stage2.log                     # where the monitor stopped / gh WARNINGs
tail -60 outputs/sft-stage2-gemma4-31b/train.log        # the actual fault
ls -lt outputs/sft-stage2-gemma4-31b/crashlogs/ | head  # did it archive this crash?
gh auth status                                          # token still valid on the host?

Reading it:

Symptom	Cause	Action
monitor not running + no new crashlogs/ entry	monitor was dead when training faulted — silent crash	relaunch (step 2)
monitor running + new crashlogs/ entry but no #120 row	crash caught but gh failed — look for WARNING: gh ... failed in monitor_stage2.log	re-auth gh / check network
train.log tail shows Memory access fault … Page not present	the known gfx1201 Tensile GEMM page fault (#113)	normal for the crawl — monitor reseeds + resumes
monitor posts ⛔ STALLED	8 consecutive no-progress retries — wedged GPU KFD	make diagnose-gfx1201-fault; full recovery in docs/HANDOFF_STAGE2_CRAWL.md

Per-crash full tracebacks + dmesg snapshots are kept in outputs/sft-stage2-gemma4-31b/crashlogs/.

[ulises-c]

Next experiments — M-sweep to crack the standalone repro (instructions)

Session 11 (full writeup in #113) found the SYNC-deterministic log records the faulting M as 607, not 608. Every standalone probe so far used M=608 — the even neighbor, from the less-reliable async capture. A second SYNC run faulted at M=439. Both are odd; the tile divergence (standalone → MT128x128x32, production → MT64x64x64) may simply be the off-by-one M crossing a Tensile size-predicate boundary.

Run on the training box (R9700), not the dev box. One GPU — do NOT run while monitor_stage2.sh is crawling. make gpu-preflight first.

1. Sweep M (PRIMARY — minutes)

make gpu-preflight
AMD_SERIALIZE_KERNEL=3 HIP_LAUNCH_BLOCKING=1 AMD_LOG_LEVEL=3 TORCH_USE_HIPBLASLT=0 \
  uv run --no-sync python scripts/probe_gfx1201_tile.py --shape-3d --ckpt \
  --m-sweep 439,512,576,606,607,608,609,640 2>&1 | tee probe_msweep.log

# which tile did each M dispatch?
grep -E 'SWEEP M=|MT[0-9]+x[0-9]+x[0-9]+.*ISA1201' probe_msweep.log \
  | grep -oE 'SWEEP M=[0-9]+|MT[0-9]+x[0-9]+x[0-9]+'

Read the result:

• Some M -> MT64x64x64 (...DTVB1...): that M is the standalone-repro shape. Go to step 2.
• All M -> MT128x128x32: M is ruled out — the trigger is the fused ..._Bias_..._UserArgs_... epilogue. Go to step 3.

2. If an M dispatched MT64x64x64 — attempt the standalone crash

Re-run the dense-VA repro (scripts/repro_gfx1201_dense.py) at that exact M with the ~20 GB filler so the wild address lands in a mapped band. A crash here = the standalone reproducer #7992 needs. Capture the ShaderName + fault address and paste both into the upstream draft, then post #7992.

3. If no M dispatched MT64x64x64 — test the epilogue path

... uv run --no-sync python scripts/probe_gfx1201_tile.py --shape-3d --ckpt --m 607 --no-bias 2>&1 | tee probe_nobias.log
... uv run --no-sync python scripts/probe_gfx1201_tile.py --shape-3d --ckpt --m 607          2>&1 | tee probe_bias.log

If MT64x64x64 appears only with bias, the trigger is the fused bias epilogue (internal Tensile dispatch, not reachable from rocblas-bench or a plain Linear4bit) — report that as kernel-confirmation-only in #7992, no standalone crash expected, and fall back to the HIP rocblas_gemm_ex + bias reproducer.

Probe changes (already committed)

scripts/probe_gfx1201_tile.py: default --m is now 607 (SYNC-confirmed), added --m-sweep <csv> (multiple M in one process, banner per M) and --no-bias. Docstring carries the run matrix + detection greps.

Upstream report corrections (#7992)

• M=608 -> M=607; state M = sample token length, data-order-dependent (439/607 observed).
• Trigger is column-major B at ldb=K=21504 specifically (same-M K=5376/16384 col-major GEMMs run clean in the same step).
• Reframe "probabilistic" -> data-order-deterministic; give the fixed data_seed=42 recipe; note seed rotation is a crawl mitigation, not a bug property.