gfx1201 (RDNA4) rocBLAS/Tensile ISA1201 GEMM page fault — upstream report reference

[ulises-c]

Summary

During Gemma‑4‑31B QLoRA training on an AMD R9700 (gfx1201 / RDNA4, ROCm 7.2), a reproducible GPU page fault ("page not present") occurs inside a rocBLAS Tensile GEMM in the QLoRA backward pass. The fault survives full kernel serialization (so it is not an async race), reproduces on a verified‑clean GPU (so it is not a dirty‑KFD cascade), and the faulting address lies in unmapped space ~1 GB from both operands of the faulting GEMM (so it is not an undersized operand descriptor or an OOM).

This issue is the internal system‑of‑record consolidating the diagnosis from PR #101. Its purpose is to seed an upstream bug report against the ISA1201 GEMM kernels, filed at ROCm/rocm-libraries — component rocBLAS (projects/rocblas/) / Tensile (shared/tensile/). (The standalone ROCm/rocBLAS and ROCm/Tensile repos are retired; rocm-libraries is the source of truth.) The ready‑to‑paste upstream draft is in the comments. See the canonical ablation tracker: docs/GFX1201_FAULT_ABLATION_LOG.md.

Related: #100 (gfx1201 training investigation), #101 (PR where this was diagnosed), #109 (gfx1201 field report).

Environment

GPU	AMD Radeon AI PRO R9700, gfx1201 / RDNA4, 32 GB
ROCm	7.2
PyTorch	2.11.0 (ROCm build)
bitsandbytes	0.49.2
Backend	rocBLAS Tensile path (TORCH_USE_HIPBLASLT=0)
Workload	Gemma‑4‑31B, NF4 QLoRA, SDPA, grad‑ckpt (use_reentrant=False), seq 1280, bs 1×16, seed 42

The fault

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

The fault is a 2 MB‑aligned host‑VA address arising during the gradient‑checkpoint forward recompute inside the backward pass. The faulting op is a rocBLAS Tensile GEMM that consumes the NF4‑dequantized weight — not the bitsandbytes dequant kernel (kDequantizeBlockwise<float> / <hip_bfloat16> succeed in every cycle).

Faulting kernel name — confirmed identical in both directions

On a verified‑clean GPU at AMD_LOG_LEVEL=3, the faulting kernel — captured in both the backward pass (runs #4/#5) and the grad‑ckpt forward recompute (run #9) — is the same gfx1201‑native Tensile GEMM, macro‑tile 64×64×64, ISA1201. Full ShaderName (run #9):

Cijk_Ailk_Bjlk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x64x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB1_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA1201_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB0_LBSPPM0_LPA32_LPB0_LPM0_LRVW8_LWPMn1_MIAV1_MIWT2_2_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB2_ONLL0_PGR2_PLR1_PKA0_SIA3_SS0_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKXCCM0_TLDS0_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS32_WG32_4_1

The same tile succeeds elsewhere in the same pass, so the fault is specific to a particular call's operands/addresses, not a categorically broken kernel. Discriminating flags that point at the B‑operand path: DTVB1 (B non‑default dtype/layout), LBSPPB0 (no B LDS prefetch), NLCB2 (B double‑unrolled), VWA2_VWB1 — consistent with the column‑major stride=(1,21504) B operand computing the wild address.

Fault‑address vs. operand descriptors (forward recompute — probe‑3)

probe‑3 (commit dad8f4a) instrumented both forward‑GEMM operands — the activation A (via MatMul4Bit.forward) and the dequantized weight B (via dequantize_4bit). The pair immediately preceding the fault:

[gemm_forward_input] shape=(1, 608, 21504) ptr=0x7f655ece8000 end=0x7f65605d8000   ← faulting GEMM A (contiguous, row-major)
[dequant_probe]      shape=(21504, 5376)   ptr=0x7f63f01a0000 end=0x7f63fde20000   ← faulting GEMM B (column-major, stride=(1,21504))
← FAULT 0x7f6459a00000

Operand	ptr	end	fault inside?
A (1,608,21504) contig	0x7f655ece8000	0x7f65605d8000	No — fault ~1.0 GB below ptr
B (21504,5376) stride=(1,21504)	0x7f63f01a0000	0x7f63fde20000	No — fault ~1.2 GB above end

A full‑log scan finds 0 operands that bracket the fault address. Both operands of the faulting GEMM are present, valid, and contiguous; the fault sits in an unmapped gap ~1 GB from either cluster, not adjacent to any buffer boundary. The column‑major (stride=(1,N)) B operand is the prime suspect for a Tensile stride‑computation bug producing the wild address.

Why it is a strong upstream report

Property	Evidence
Reproducible	Faults across multiple runs / resume points (runs 4–7 in the ablation log).
Serialized — not a race	Persists under AMD_SERIALIZE_KERNEL=3 HIP_LAUNCH_BLOCKING=1 (run #3). Serialization kills races, not this fault.
Both directions, same tile	Backward fault (probe‑1) and forward‑recompute fault (probe‑3) resolve to the identical MT64x64x64 … ISA1201 ShaderName (run #9).
Not bitsandbytes	NF4 dequant kernels succeed every cycle; the fault is in the GEMM consuming the dequant output (runs #5–7).
Not an undersized descriptor	Fault is ~1 GB from any operand boundary, not adjacent to end (probe‑1, probe‑3).
Not a dirty‑KFD cascade	Reproduces on a make gpu-preflight‑verified‑clean GPU (run #5).
Not OOM	VRAM at fault is within headroom (alloc ~21 GB / reserved ~28 GB of 32 GB).

What would go in the upstream (rocm-libraries) report

• Arch: gfx1201 (RDNA4 / Navi 48), ROCm 7.2, torch 2.11.0+rocm7.2.
• Faulting kernel (forward recompute and backward — confirmed identical): full ShaderName above (…_MT64x64x64_…_ISA1201_…).
• Operand descriptors at fault (forward GEMM): A (1,608,21504) row‑major contiguous; B (21504,5376) column‑major stride=(1,21504); bf16.
• Fault address 0x7f6459a00000 lies outside both operand ranges (≥1 GB gap).
• Reproducible and serialized; gfx1201 only.

Resolved since first filing

• Forward tile name — CONFIRMED (run docs: rename resources/ to references/, add HuggingFace notes #9, AMD_LOG_LEVEL=3). The forward‑recompute faulting GEMM resolves to the same ShaderName as the backward kernel (see above). The earlier concern that they might be different tiles (the forward problem size (1,608,21504)@(21504,5376) differs) is settled by evidence: both dispatch the identical MT64x64x64 … ISA1201 tile. The upstream report can name the kernel directly for both directions.

Open / unconfirmed

• Output buffer not instrumented. A page‑not‑present is formally consistent with an OOB write to the output rather than a wild read. Deemed implausible — the output (1,608,5376) bf16 is ~6 MB and the fault is ~1 GB from any logged buffer — but it was not ruled out by instrumentation. This is the sole residual ambiguity in the wild‑read characterization.

Mitigations under evaluation (PR #101)

Tracked in docs/GFX1201_FAULT_ABLATION_LOG.md, scored on two axes (does it eliminate the fault / does it improve steps‑per‑resume):

• TORCH_USE_HIPBLASLT=1 (Arm HLT1) — tried, not a fix (run feat: Ollama remote-access setup for M4 Mac Mini consultant #11). hipBLASLt has no kernel for this MT64x64x64 DTVB1 shape (logs Cannot find the function ×6) and falls back to the same Tensile kernel; fault persists. No userspace routing around the bug exists in ROCm 7.2.
• expandable_segments:True (Arm D) — dead arm (run feat: plug-and-play WAVE HPC setup script #10): silently ignored on gfx1201/ROCm, so it only tested "no GC threshold." Fault persisted regardless → allocator placement cannot mask it.
• use_reentrant=True grad‑ckpt — de‑prioritized: targets the activation‑lifetime hypothesis (Bucket feat: add poetry run test and poetry run lint commands #1), which probe‑3 refuted (both operands valid; wild address ⇒ Bucket feat: add GitHub Actions CI with ruff linting and pytest #2).
• HSA_ENABLE_SDMA=0 — cheap, stackable; folded into the pending production steps‑per‑resume measurement.

Both routing/placement arms are now exhausted. Remaining paths are the brute‑force checkpoint crawl (viability unmeasured) or relocating the SFT run to NVIDIA (see docs/SFT_NVIDIA_MIGRATION.md); the upstream fix is the only thing that unblocks gfx1201 itself.

References

• PR fix(train): gfx1201 ROCm fixes + unsloth bnb-4bit Gemma 4 31B QLoRA #101 — full diagnostic thread (probe‑1 / probe‑2 / probe‑3, consultant reviews).
• docs/GFX1201_FAULT_ABLATION_LOG.md — canonical run/ablation tracker.
• docs/GFX1201_RDNA4_TRAINING.md §6.1 — original narrative (predates localization; the tracker supersedes its bnb‑side framing).
• Probe commits: 81cebb4 (probe‑1), 1b752c2 (probe‑2), dad8f4a (probe‑3). Run docs: rename resources/ to references/, add HuggingFace notes #9 (kernel-name.log, AMD_LOG_LEVEL=3) confirmed the forward ShaderName; run feat: plug-and-play WAVE HPC setup script #10 (11c13b6) closed Arm D.

[ulises-c]

📋 Upstream issue draft — ready to paste to ROCm/rocm-libraries

Component: rocBLAS (projects/rocblas/) / Tensile (shared/tensile/). The block below is the standalone upstream report. Environment + versions are now filled in; the one remaining artifact is the rocblas-bench reproducer line (ROCBLAS_LAYER=2 capture) — the procedure to capture it (and confirm the column-major-B trigger in the same run) is in the Reproducer section.

---

Title: gfx1201 (RDNA4): ISA1201 Tensile GEMM MT64x64x64 computes out-of-bounds address (GPU page fault) on a column-major B operand

---

Summary

On gfx1201 (RDNA4) under ROCm 7.2.3, a Tensile-generated GEMM kernel (Cijk_…_MT64x64x64_…_ISA1201) triggers a reproducible GPU page fault ("page not present"). The faulting access lands at a host-VA address ~1 GB outside the bounds of both GEMM operands, indicating the kernel computes an out-of-bounds address rather than reading a valid-but-unexpected buffer. The fault is deterministic enough to reproduce on a clean GPU and persists under full kernel serialization (AMD_SERIALIZE_KERNEL=3 HIP_LAUNCH_BLOCKING=1), so it is not an asynchronous race. The same kernel and tile succeed on other problem sizes in the same workload; the fault is specific to a call whose B operand is column-major (stride=(1, K)).

The workload is Gemma-4-31B NF4 QLoRA fine-tuning (PyTorch + bitsandbytes); the fault is in the rocBLAS GEMM consuming the dequantized weight, not in the bitsandbytes dequant kernel (verified — see "Ruled out").

Environment

GPU	AMD Radeon AI PRO R9700, gfx1201 / RDNA4 (Navi 48), 32 GB; VBIOS 113-1E4990U-S83, SKU 1E4990U, rev 0xc0
ROCm / HIP	7.2.53211-364a905 (ROCm 7.2.3, HIP 7.2.26015)
rocBLAS	7.2.3-1 (Tensile bundled; no separate package)
hipBLASLt	7.2.3-1
OS / kernel	CachyOS, kernel 6.x (7.0.10-2-cachyos), in-kernel (inbox) amdgpu driver (not amdgpu-dkms), linux-firmware-amdgpu 20260519
PyTorch	2.11.0+rocm7.2 (torch.version.hip = 7.2.26015)
bitsandbytes	0.49.2
Backend	rocBLAS Tensile path (TORCH_USE_HIPBLASLT=0; see "Ruled out" — hipBLASLt has no kernel for this shape and falls back to Tensile regardless)

The correct-arch Tensile library is present and loaded: /opt/rocm/lib/rocblas/library/ contains both TensileLibrary_lazy_gfx1201.dat and Kernels.so-000-gfx1201.hsaco, and the faulting kernel is ISA1201 (gfx1201's ISA). This rules out the wrong-library-lookup failure mode of #7192 — we are running the genuine gfx1201 kernel, which itself faults.

Faulting kernel (full ShaderName)

Captured at AMD_LOG_LEVEL=3 on a clean GPU. The same kernel variant faults in both the forward and backward GEMM of the workload:

Cijk_Ailk_Bjlk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x64x64_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB1_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA1201_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB0_LBSPPM0_LPA32_LPB0_LPM0_LRVW8_LWPMn1_MIAV1_MIWT2_2_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA1_NLCB2_ONLL0_PGR2_PLR1_PKA0_SIA3_SS0_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKXCCM0_TLDS0_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS32_WG32_4_1

Tile: MT64x64x64, MI16x16x1, ISA1201. Flags distinguishing the B-operand path: DTVB1 (B non-default dtype/layout), LBSPPB0 (no B LDS prefetch), NLCB2 (B double-unrolled), VWA2_VWB1.

The faulting call — operands vs. fault address

GEMM dimensions (derived from the logged operands): M=608, N=5376, K=21504, bf16, batch 1. A is row-major contiguous (M×K); B is column-major (K×N), stride=(1, 21504); C is (M×N) bf16 (~6 MB). The pair logged immediately before the fault:

A (gemm input)   shape=(1, 608, 21504)  stride=(13074432, 21504, 1)  contig=True   ptr=0x7f655ece8000  end=0x7f65605d8000
B (dequant wt)   shape=(21504, 5376)    stride=(1, 21504)            contig=False  ptr=0x7f63f01a0000  end=0x7f63fde20000
FAULT: 0x7f6459a00000

Operand	range	fault inside?
A (1,608,21504) row-major contig	0x7f655ece8000 … 0x7f65605d8000	No — fault ~1.0 GB below ptr
B (21504,5376) col-major stride=(1,21504)	0x7f63f01a0000 … 0x7f63fde20000	No — fault ~1.2 GB above end

A full scan of the run logs 0 GEMM operands that bracket 0x7f6459a00000. The fault sits in unmapped space ~1 GB from either operand cluster — not adjacent to any buffer boundary, which rules out a simple off-by-N or undersized-descriptor read.

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

(Address varies run to run — e.g. 0x7f0ac2e00000, 0x7f29eb600000 in other runs — but is consistently ~1 GB from the operands and 2 MB-aligned.)

Reproducibility

• Deterministic on a clean GPU, faults reproducibly within a few steps of resuming the workload.
• Persists under serialization — AMD_SERIALIZE_KERNEL=3 HIP_LAUNCH_BLOCKING=1 → not an async race.
• Both directions — forward-recompute GEMM and backward GEMM dispatch the identical tile and both fault.

Minimal reproducer

TODO — rocblas-bench line (last artifact, capture in progress). rocblas-bench is installed (/opt/rocm/bin/rocblas-bench). Capturing one ROCBLAS_LAYER=2 run emits the exact rocblas-bench -f gemm_ex --transposeA … --transposeB … -m 608 -n 5376 -k 21504 --lda … --ldb … --ldc … line for the faulting GEMM; running that line directly reproduces the fault independent of PyTorch/bitsandbytes, isolating it entirely within rocBLAS/Tensile. Will be attached here as a follow-up.

Until then the trigger is reachable via any gemm_ex with this selection: bf16, M=608/N=5376/K=21504, transposed/column-major B (the DTVB1 path), on gfx1201 with the Tensile backend.

Ruled out (to save triage time)

Hypothesis	Evidence
Asynchronous / concurrency race	Persists under AMD_SERIALIZE_KERNEL=3 HIP_LAUNCH_BLOCKING=1.
Caller passed an undersized / wrong-shape operand	Both operands logged, valid, contiguous as expected; fault is ~1 GB from either boundary, not adjacent.
bitsandbytes NF4 dequant kernel	Dequant kernels succeed every iteration; fault is in the GEMM consuming their output.
hipBLASLt	TORCH_USE_HIPBLASLT=1 was tried — hipBLASLt has no kernel for this MT64x64x64 DTVB1 shape (logs Cannot find the function ×6) and falls back to this same Tensile kernel; fault persists.
Allocator placement / fragmentation	PYTORCH_HIP_ALLOC_CONF GC-threshold and expandable_segments changes do not move or eliminate the fault (expandable_segments is in fact ignored on gfx1201).
OOM	VRAM at fault ~21 GB allocated / ~28 GB reserved of 32 GB — within headroom.

Open / unconfirmed

• Column-major B as the trigger is observed, not yet proven. Every faulting call has a stride=(1,K) B operand and the DTVB1 kernel flag, and the same tile succeeds on other shapes — strong correlation. Forcing the dequantized weight row-major at the framework (Python) level does not change the outcome and does not test the hypothesis: bitsandbytes sets the transpose via the BLAS transB flag (the A @ Wᵀ call structure), not via the tensor's physical stride, so a Python .contiguous() never reaches the rocBLAS descriptor (the DTVB1 kernel still dispatches and still faults). The definitive test is at the rocBLAS layer: run the captured rocblas-bench line with --transposeB T (reproduces) vs. an equivalent --transposeB N (row-major B → DTVB0, a different tile) and compare. That experiment is folded into the reproducer capture above.
• Output buffer not instrumented. A page-not-present is formally consistent with an OOB write to C rather than a wild read of an input. Considered implausible (C is ~6 MB and the fault is ~1 GB away) but not ruled out by instrumentation.

Impact

Blocks NF4 QLoRA fine-tuning of large models (Gemma-4-31B) on gfx1201/RDNA4 under ROCm 7.2 — any backward pass that dispatches this tile faults. No userspace workaround found (hipBLASLt lacks the kernel; allocator/serialization knobs don't help).

Possibly related

• rocm-libraries#6166 — rocBLAS unit tests (dot_ex) on gfx1201 produce the identical Memory access fault … Page not present signature with no PyTorch/bitsandbytes in the picture. Strongest corroboration that gfx1201 rocBLAS kernels page-fault at the library level; likely an in-house-runnable repro of the same class of bug.
• rocm-libraries#4097 — same kernel family (MT64x64x64 … DTVB1 … ISA1201) on Windows, but the missing-kernel (hipErrorNotFound) mode rather than a present-but-faulting kernel. Suggests the DTVB1 ISA1201 MT64x64x64 tiles are the trouble spot for gfx1201.
• rocm-libraries#7192 — same GPU (R9700); a wrong-Tensile-file (gfx1200.dat) lookup. Ruled out as our mechanism (we confirmed the gfx1201 library is present and the kernel is ISA1201), included as same-hardware context.

[ulises-c]

TODO — strengthen the upstream report (ROCm/rocm-libraries#7992)

The diagnosis in #7992 is solid; the one gap it flags itself is a standalone reproducer (no PyTorch/bitsandbytes). That single artifact is what lets a maintainer act in minutes instead of trusting the write-up. Priority order:

1. Standalone reproducer (make-or-break)

• Don't hand-write rocblas-bench args — the transpose/lda/ldb encoding for the col-major stride=(1,21504) B is easy to get wrong, and a wrong transpose flag silently selects a different kernel that won't fault. Capture the real invocation from a live faulting run instead.
• Set rocBLAS bench logging on a make gpu-preflight'd resume: ROCBLAS_LAYER=2 (bitmask: 1=trace, 2=bench, 4=profile) + ROCBLAS_LOG_BENCH_PATH=<file>. Verify the bitmask value on-box first — a wrong value just yields an empty log.
• No need to capture the faulting call itself (it never returns → may never flush). The fault is shape-deterministic (sticky at step ~21). Grep the bench log for any gemm_ex line with M=608 N=5376 K=21504 bf16 → that's the exact arg encoding to replay.
• Success criterion is not "it runs" — it's "it dispatches the same ShaderName AND faults." Replay the captured bench line under AMD_LOG_LEVEL=3 and confirm it dispatches …MT64x64x64…ISA1201…DTVB1….
  • Faults w/ same ShaderName → paste the one-line repro into #7992. Done.
  • Does not fault → that's a finding: the bug needs the fused/UserArgs path. Report the negative + fall back to a ~40-line HIP program calling rocblas_gemm_ex with bias on the same entry point.

2. Shape minimization (strong second — only after a standalone repro exists)

• Shrink M/N/K until the fault disappears. The boundary pinpoints the stride/offset arithmetic that overflows — that is the bug. Can't be done via training runs (crash cycles are minutes + data-order sticky); only works on the standalone repro.

3. Cross-link related gfx1201 memory-fault issues

• Ping / cross-link #6166 (dot_ex unit-test page fault, same arch, OPEN, in triage, assigned zichguan-amd, has a Docker repro). Different kernel but same gfx1201 fault class → routes #7992 to the same owner. (#4097 / #7192 already linked.)

4. Polish (one line each — don't let it crowd out the repro)

• Exact package versions: dpkg -l | grep rocblas (or rpm -qa) for precise rocblas/tensile/hipblaslt versions + stock-vs-dev-build. "ROCm 7.2.3" alone leaves the Tensile library version ambiguous.
• Close the "OOB-write vs wild-read" residual ambiguity by folding output-buffer (C) logging into the next instrumented run — don't headline it; the report already deems a write implausible.
• Skip Tensile solution-index capture — the full ShaderName is already better.

Artifacts to author for the R9700 (dev box is NVIDIA, can't run/verify there): (1) the bench-capture + same-ShaderName-assert replay script, (2) the HIP rocblas_gemm_ex+bias fallback repro.

[ulises-c]

Scripts ready — bench capture + replay

Two scripts added in scripts/ to produce the standalone reproducer (TODO item 1 from the previous comment).

scripts/capture-rocblas-bench.sh

Sets ROCBLAS_LAYER=2 ROCBLAS_LOG_BENCH_PATH=<file> on a single one-shot resume, runs until the GPU page fault (~21 steps after the checkpoint), then greps the bench log for any gemm_ex call at M=608 N=5376 K=21504 bf16 and prints the candidate line(s).

Key design points:

• Verifies ROCBLAS_LAYER=2 produces output before kicking off the 19 GB model load (probe call on a tiny GEMM first)
• Does make gpu-preflight before the run
• Does NOT hand-write the transpose/lda/ldb args — captures them from the live Tensile dispatch to avoid the DTVB1 encoding mistake
• Bounded (TRAIN_MAX_STEPS = latest_ckpt + 50) so it doesn't run indefinitely

Do not run while the production crawl (monitor_stage2.sh) is active — one GPU.

scripts/replay-rocblas-bench.sh "<line>"

Takes the captured bench line and replays it under AMD_LOG_LEVEL=3 AMD_SERIALIZE_KERNEL=3, then checks:

1. Does it produce a GPU page fault? → standalone repro confirmed
2. Does the ShaderName match ...MT64x64x64...ISA1201...DTVB1...? → correct tile confirmed

If the replay does not fault, it means the bug requires the UserArgs/bias-fused path (consistent with the kernel name BH_Bias_HA_S_SAV_UserArgs). That case falls back to the HIP rocblas_gemm_ex + bias reproducer (still to author).

Package version note (polish TODO)

pacman -Q confirms the installed packages match what's already in the upstream draft:

• rocblas 7.2.3-1 (stock Arch package, not a dev build)
• hipblaslt 7.2.3-1

No change needed in the environment table.

Cross-link ping: rocm-libraries#6166

#6166 (dot_ex unit-test page fault, same Memory access fault … Page not present signature on gfx1201, assigned zichguan-amd, has a Docker repro) is the strongest corroboration that gfx1201 rocBLAS kernels fault at the library level with no PyTorch/bitsandbytes in the picture. The upstream report already lists it in "Possibly related"; worth pinging that issue directly (or leaving a cross-reference comment there) once the standalone bench repro is confirmed, so it routes to the same owner.

Remaining TODO

• Run bash scripts/capture-rocblas-bench.sh after the current crawl run crashes to get the bench line
• Run bash scripts/replay-rocblas-bench.sh "<line>" to confirm fault + ShaderName
• If replay confirms → paste one-line repro into the upstream draft (comment 1 above) and post to ROCm/rocm-libraries
• If replay does NOT fault → author HIP rocblas_gemm_ex + bias fallback reproducer
• Optional: add C output-buffer (C tensor) logging to the probe to close the "OOB-write vs wild-read" residual ambiguity

[ulises-c]

Session findings — bench capture + replay runs (2026-06-05)

What we ran

Three scripts against a freshly verified-clean GPU (post make gpu-preflight):

1. scripts/collect-gfx1201-report-env.sh — env snapshot
2. scripts/capture-rocblas-bench.sh — one-shot training resume with ROCBLAS_LAYER=2 ROCBLAS_LOG_BENCH_PATH=<file> to ~50 steps past checkpoint-1230
3. scripts/replay-rocblas-bench.sh — two manual bench-line attempts (both B-operand encodings)

Pre-req fix: the venv had regressed to torch+cu130 (CUDA build, GPU not visible). Fixed with make install → torch==2.11.0+rocm7.2, hip: 7.2.26015, GPU visible.

---

Finding 1 — ROCBLAS_LAYER=2 does not log through PyTorch's hipBLAS path

The probe (direct rocblas-bench -f gemm -m 4 -n 4 -k 4 call) produced output — ROCBLAS_LOG_BENCH_PATH is recognized by librocblas.so. But the training run produced zero bytes in the bench log despite faulting correctly.

Root cause: PyTorch calls libhipblas.so → internal rocBLAS dispatch, not the public rocblas_gemm_ex C API entry points that ROCBLAS_LAYER hooks. The logging is bypassed entirely by the hipBLAS wrapper path. This was confirmed by also testing ROCBLAS_LAYER=1 (trace) — also silent through PyTorch.

ROCBLAS_LAYER bench/trace logging only fires on direct calls to the rocBLAS C API, not on calls routed through hipBLAS. This is a hard limitation of the logging architecture.

---

Finding 2 — Plain rocblas-bench gemm_ex always selects ISA000 (MLIR generic), never ISA1201

Tried both B-operand encodings for M=608, N=5376, K=21504, bf16:

Command	Kernel dispatched	Faulted?
--transposeA T --transposeB N --lda 608 --ldb 21504 --ldc 608	ISA000 MT64x32x8 … DTVB0	No
--transposeA T --transposeB T --lda 21504 --ldb 5376 --ldc 608	ISA000 MT32x32x16 … DTVB0	No

Both encodings route to the MLIR-compiled generic (ISA000) library, not the gfx1201-specific Tensile (ISA1201) library. The DTVB0 flag (default B dtype/layout) in both replays confirms neither encoding matches the faulting call's DTVB1 path.

Checked: rocblas_gemm_ex3 / rocblas_epilogue_t are not present in the ROCm 7.2.3 rocblas.h — the UserArgs/Bias path is purely internal Tensile dispatch, not exposed in the public C API headers, and not reachable from rocblas-bench.

Conclusion: The ISA1201 MT64x64x64 … DTVB1 … BH_Bias_HA_S_SAV_UserArgs kernel is only selected through the bitsandbytes matmul_4bit → dequantize_4bit → rocBLAS call chain. No userspace routing around the training stack exists to trigger it directly.

---

Finding 3 — Fault confirmed again + new artifact: VRAM allocation failure

The capture run (clean GPU, AMD_LOG_LEVEL=1) produced the same fault:

Memory access fault by GPU node-1 (Agent handle: 0x556a2bfc58f0) on address 0x7f9c14800000.
Reason: Page not present or supervisor privilege.

New: AMD_LOG_LEVEL=1 surfaced a VRAM allocation failure ~6 minutes before the fault that was invisible in prior runs (which used the default log level):

:1:rocdevice.cpp  :2189: Fail allocation local memory
:1:rocdevice.cpp  :1877: Failed creating memory
:1:memory.cpp     :357:  Video memory allocation failed!
:1:memory.cpp     :317:  Can't allocate memory size - 0x38E00000 bytes! (~936 MB)
:1:rocdevice.cpp  :2256: failed to create a svm hidden buffer!

Timestamps: allocation failure at 121491877606 µs, fault at 121870395584 µs (~378 s gap, corresponding to several training steps). The gap makes a direct causal link unlikely (training continued for minutes after the allocation failure before faulting), but it represents a new observable in the picture.

Also visible in the same log — the HIP module cache misses for the exact faulting kernel before the fault:

:1:hip_code_object.cpp  :119: Cannot find the function: Cijk_Ailk_Bjlk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x64x64_…_ISA1201_…_DTVB1_…

(Multiple modules searched before the kernel is loaded from .hsaco; this is normal lazy-load cache miss, not an error. The kernel still runs and faults.)

---

Updated TODO state

• ROCBLAS_LAYER=2 bench capture — dead end: hipBLAS path bypasses all ROCBLAS_LAYER logging
• Plain rocblas-bench gemm_ex replay — dead end: always dispatches ISA000, never ISA1201
• Minimal Python reproducer — next and only remaining path to a standalone repro. A ~20-line script using bitsandbytes.nn.Linear4bit at the exact faulting shape (in_features=21504, out_features=5376, batch×seq=608) exercises the exact matmul_4bit → dequant → UserArgs GEMM chain without loading a 19 GB model. No trainer, no data. If it faults, it's the standalone repro.

[ulises-c]

Session 8 findings — SYNC log parsed + improved reproducer written

SYNC diagnostic log results

docs/diagnostics/diag-sync-probe-step1230-20260605-124008.log (26 GB, 230M lines) has been fully parsed. Key findings:

Fault confirmed:

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

• 0x7efb33000000 is 2 MB-aligned ✓ — consistent with all prior crashes
• Fault occurred within ~1–2 steps of resuming checkpoint-1230 (SYNC mode = deterministic)

Same faulting kernel confirmed at crash point:

ShaderName: Cijk_Ailk_Bjlk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x64x64_MI16x16x1_…_ISA1201_…_DTVB1_…_VWA2_VWB1_…

Identical to runs #5, #7, #9 — confirms the kernel hasn't changed across resumes.

Last B-operand probe immediately before fault (via BNB_DEQUANT_PROBE=1):

[dequant_probe] shape=(21504, 5376) stride=(1, 21504) contig=False dtype=torch.bfloat16
  ptr=0x7ef978800000  end=0x7ef986480000

B is col-major (stride=(1,21504) = stride=(1,K)), non-contiguous — the DTVB1 path.
Fault is 6.7 GB above B.end (2 MB-aligned). ASLR shifts the absolute address each run
but the 2 MB-alignment is invariant across all 13+ observed faults.

---

Improved standalone reproducer

scripts/repro_gfx1201_dense.py has been written. It addresses the known failure mode of
repro_gfx1201_rocblas.py (sparse VA → wild address lands outside all VMAs → no observable page fault).

The fix: pre-allocate ~20 GB of GPU tensors before building the Linear4bit layer.
The filler occupies a large VA band. The layer's B matrix is placed at a higher VA; the
wild address (B.ptr + ~6.7 GB) is more likely to fall inside one of the filler chunks (or in
the gap immediately above them), which the HMM page-walk resolves as page-not-present → crash.

Run:

make gpu-preflight   # GPU must be clean
AMD_SERIALIZE_KERNEL=3 HIP_LAUNCH_BLOCKING=1 AMD_LOG_LEVEL=3 TORCH_USE_HIPBLASLT=0 \
  uv run --no-sync python scripts/repro_gfx1201_dense.py 2>&1 | tee repro_dense.log

# Confirm kernel dispatched regardless of crash:
grep -o 'MT64x64x64[^[:space:]]*ISA1201[^[:space:]]*' repro_dense.log | head -3

Success condition: process killed with Memory access fault … Page not present.
Fallback: if still no crash, the kernel dispatch is confirmed by the ShaderName grep
and the bug is real but requires the full model VA footprint to observe — report that
finding upstream as well.

---

Updated TODO for upstream #7992

• rocblas-bench capture — confirmed dead end (ROCBLAS_LAYER=2 silent through hipBLAS; rocblas-bench always dispatches ISA000, never ISA1201 for this shape)
• SYNC log parsed — fault + ShaderName confirmed, new B-operand addresses logged
• Run repro_gfx1201_dense.py — test whether 20 GB VA filler is sufficient to trigger the fault without the full model
• Post update to upstream #7992:
  • Close the rocblas-bench TODO (dead end, explain why)
  • Share SYNC confirmation + new B-operand addresses
  • Attach dense reproducer result (crash or kernel-dispatch-only)

---

What AMD sees right now in #7992

The last comment there still says "I'll run some tests when training is done" (from when training was actively running). The issue body still has the rocblas-bench reproducer as "in progress." Nick Kuo (assigned) is waiting on that. Next session: post the update and the dense reproducer result.

[ulises-c]

Session 9 — dense-VA reproducer result + tile-selection finding

Run summary

GPU preflight: PASS (VRAM idle 57 MB, fwd+bwd OK)
repro_gfx1201_dense.py — 2026-06-06
  torch:        2.11.0+rocm7.2
  bitsandbytes: 0.49.2
  Device:       AMD Radeon AI PRO R9700  (31.9 GB)
  ENV:          AMD_SERIALIZE_KERNEL=3  HIP_LAUNCH_BLOCKING=1  AMD_LOG_LEVEL=3  TORCH_USE_HIPBLASLT=0

Pre-allocating VA filler: target 20 GB in 512 MB chunks …
  Filler: 40 chunks, 20.0 GB allocated

Building Linear4bit layer  (in=21504, out=5376, bias=True, nf4, bf16) …
After warmup: allocated=20644 MB  reserved=21442 MB

Running 200 forward+backward passes under SYNC …
  iter   20/200 — OK
  iter   40/200 — OK
  …
  iter  200/200 — OK

Completed 200 iterations without a catchable exception.

Result: no crash.

---

Kernel-dispatch finding — tile selection mismatch

The MT64x64x64 tile was NOT dispatched in the standalone reproducer. Instead, ISA1201 DTVB1 dispatched MT128x128x32 — 401 times, all successful:

ShaderName: Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x128x32_MI16x16x1_SN_…_ISA1201_…_DTVB1_…_VWA1_VWB1_…

The production faulting kernel was:

ShaderName: Cijk_Ailk_Bjlk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT64x64x64_MI16x16x1_SN_…_ISA1201_…_DTVB1_…_VWA2_VWB1_…

Key difference: Alik_Bljk (standalone) vs Ailk_Bjlk (production). The index ordering of A changes — this is a GEMM transpose-flag difference. The production fault is specific to the Ailk_Bjlk layout (column-major A / row-major B), which only arises during the grad-checkpoint recompute path, not a plain layer(x) forward call.

The standalone script calls layer(x) directly in forward mode; grad-ckpt recompute sets up the GEMM from inside torch.utils.checkpoint with a different activation layout. That layout change routes Tensile to MT64x64x64 instead of MT128x128x32, and that tile is the one with the wild-address bug.

---

What this means for the upstream report

• ✅ ISA1201 DTVB1 path is exercised by bitsandbytes.nn.Linear4bit at this shape
• ✅ MT128x128x32_ISA1201_DTVB1 runs successfully — that tile is not broken
• ❌ MT64x64x64_ISA1201_DTVB1 cannot be triggered without the grad-checkpoint recompute call path (or an equivalent Ailk_Bjlk-layout GEMM call)
• The dense-VA approach solved the VA footprint problem; the remaining gap is the call-site transpose-flag context

---

Updated TODO for upstream #7992

• rocblas-bench capture — dead end (rocblas-bench always dispatches ISA000)
• SYNC log parsed — fault + ShaderName confirmed, B-operand addresses logged
• Dense-VA reproducer run — no crash; MT128x128x32 dispatched, MT64x64x64 not triggered
• Post update to upstream #7992:
  • Share dense-VA result + tile-selection finding
  • Clarify that a standalone repro must replicate the Ailk_Bjlk (col-major A) call, not just the shape
  • The upstream report body can be updated to note this as the only remaining reproducibility gap

[ulises-c]

Session 9 — tile-selection probe grid: all variants dispatch MT128x128x32

What was tried (9 probe variants, all on clean GPU after reboot)

New script: scripts/probe_gfx1201_tile.py — runs 3 fwd+bwd iters with AMD_LOG_LEVEL=3,
checks which ShaderName is dispatched. No VA filler (pure kernel-name probe).

Config	Kernel dispatched
2D (608,21504) plain (baseline from dense run)	MT128x128x32
3D (1,608,21504) plain	MT128x128x32
3D (1,608,21504) + ckpt(reentrant=False)	MT128x128x32
2D + ckpt(reentrant=False)	MT128x128x32
3D + ckpt(reentrant=True)	MT128x128x32
2D + torch.autocast	MT128x128x32
2D + no warmup (cold Tensile cache)	MT128x128x32
2D + non-contiguous A (stride=[43008,1])	MT128x128x32

Exhausted: shape (2D/3D), checkpoint mode (plain/reentrant_false/reentrant_true), autocast, cold cache, non-contiguous input. None select MT64x64x64.

---

What the bitsandbytes forward actually calls

Traced F.linear in MatMul4Bit.forward:

• dequantize_4bit(B, qs) returns (21504, 5376) col-major, stride=(1, 21504) ← this matches production probe-3 exactly
• .t() → (5376, 21504) row-major, contiguous
• F.linear(A, W_t, bias) calls addmm(bias, A, W_t.t()) where W_t.t() = (21504, 5376) col-major → this IS the DTVB1 path

Both standalone and production make structurally identical BLAS calls:

• A: (M, 21504) row-major
• B: (21504, 5376) col-major, stride=(1, 21504) — DTVB1

Yet production selects MT64x64x64 Ailk_Bjlk and the standalone always selects MT128x128x32 Alik_Bljk.

---

Leading hypothesis: system rocBLAS vs torch-bundled library

The torch-bundled rocBLAS at torch/lib/rocblas/library/TensileLibrary_lazy_gfx1201.dat appears to lack MT64x64x64 kernel entries for gfx1201 — no MT64x64x64 strings found. The system ROCm 7.2 library at /opt/rocm/lib/rocblas/library/ has Ailk_Bjlk kernel files and its TensileLibrary_lazy_gfx1201.dat may include MT64x64x64.

If the production training process loaded the system library (e.g., if /opt/rocm/lib appeared in LD_LIBRARY_PATH before the venv, or if ROCBLAS_TENSILE_LIBPATH was set), Tensile would select from a different kernel menu — and MT64x64x64 could be the preferred tile for this problem.

Next step (blocked by Railguard path fence — needs approval):

# Check if system library has MT64x64x64:
strings /opt/rocm/lib/rocblas/library/TensileLibrary_lazy_gfx1201.dat \
  | grep -o 'MT[0-9]*x[0-9]*x[0-9]*' | sort | uniq -c | sort -rn | head -20

# If yes, probe with system library:
ROCBLAS_TENSILE_LIBPATH=/opt/rocm/lib/rocblas/library \
  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 2>&1 | tee probe_system_rocblas.log

If MT64x64x64 appears → add 20 GB filler and attempt crash.
If not → tile is unreachable from standalone; report upstream as kernel-only confirmation (no standalone crash).

---

Immediate TODO for upstream #7992

Regardless of whether the crash is reproduced standalone, the upstream report can now include:

• ✅ DTVB1 ISA1201 path confirmed exercised by bitsandbytes.nn.Linear4bit at this shape
• ✅ MT128x128x32 runs 200+ fwd+bwd iters without fault (tile is NOT broken)
• ❌ MT64x64x64 not dispatched from standalone — tile selection divergence unexplained (library hypothesis pending)
• The bitsandbytes BLAS call is fully traced; the operand layout is the same as production

[ulises-c]

Session 10 — system-rocBLAS hypothesis disproved

System library vs. torch-bundled library — result

Compared probe_3d_plain.log (torch-bundled library) against probe_system_rocblas.log (ROCBLAS_TENSILE_LIBPATH=/opt/rocm/lib/rocblas/library). Both dispatch identically:

MT128x128x32 × 18   (all ISA1201 DTVB1)
MT64x64x64   × 0

The two apparent MT64x64x64 hits in the probe logs were the probe script's own print() statements, not actual kernel dispatches. The system library and torch-bundled library are equivalent for gfx1201 — the hypothesis is false.

Updated standalone reproducer status

Approach	Result
9 probe variants (2D/3D, ckpt, autocast, cold cache, noncontig A)	All → MT128x128x32
Dense-VA (20 GB filler + 200 fwd/bwd iters)	No crash; MT128x128x32 × 401
System library override (ROCBLAS_TENSILE_LIBPATH)	Same as torch-bundled — MT128x128x32

Tile selection divergence (standalone→MT128x128x32, production→MT64x64x64) is not explained by library path. The trigger must lie elsewhere in the production call site context (model VA footprint, multi-layer Tensile cache state, or a subtle GEMM-descriptor difference not captured by shape alone).

Training

Monitor running. Step 1305/4826 (27%) as of 2026-06-06 19:35 PST. One crash this session, recovered automatically.

Next

Post all findings to upstream ROCm/rocm-libraries#7992 — no further standalone investigation pending.