P0: IQ3_S/IQ2_S GEMV kernels fail to load on cc>=12 (GB10/Blackwell) — GH-480 PTX patch injects .reg inside a global array initializer

P0: IQ3_S / IQ2_S GEMV kernels cannot load on compute capability ≥ 12 (GB10 sm_121, RTX 50xx sm_120)

The GH-480 backward-branch pass (patch_backward_branches_sm121) injects its predicate declaration inside a global array initializer, so ptxas rejects the module. Root cause pinned by aprender-dd, from the #3968 conformance harness; verified by aprender-83 on the release cut.

Reproduced on the 0.69.1 release binary (X2)

Binary: apr 0.69.1 (d8a6df53a) on gx10 (GB10, sm_121), sha256 53e506d466bb0262…0632.
Model: Qwen3.5-0.8B-UD-IQ2_XXS, which holds 5 IQ3_S + 5 IQ2_S tensors.

apr run Qwen3.5-0.8B-UD-IQ2_XXS.gguf --prompt "What is 2+2?" --max-tokens 32 --temperature 0 --gpu --format json
[PTX-JIT] Try 1 failed: .visible .entry iq3_s_gemv_warp_reduce(, target: sm_121, PTX: 11097 bytes, result: 218
[PTX-JIT] Error log: ptxas application ptx input, line 8; fatal   : Parsing error near '.reg': syntax error
warning: GPU forward failed ... 'qwen35_cuda_deltanet' not supported: CUDA module loading failed: try1=218 try2=218 (JIT target: sm_121) — rejecting the CUDA path and falling back (fail-closed)
error: Backend unavailable: gpu was forced and selected, but its runtime attempt on this model failed and the generation ran on CPU. Refusing to report that as success
rc=14   provenance: {"requested": "gpu", "ran": "cpu", "fell_back": true}

Behaviour: fail-closed, never a wrong answer.

Under --gpu apr exits 14 and names the reason.
Without --gpu, the provenance says fell_back: true and generation runs on the CPU.
A model that holds IQ3_S or IQ2_S therefore has no CUDA path on GB10 or Blackwell.

Root cause (`crates/aprender-gpu/src/driver/ptx_patch.rs`, at `d8a6df5`)

emit_patched_line sets in_body on the first line whose trimmed text is { or ends with { (ptx_patch.rs:102).
In IQ3_S that first line is the codebook initializer .global .align 4 .u32 iq3s_grid_g[512] = {. IQ2_S builds a table the same way.
The pass then writes .reg .pred %p_jw; + setp.ne.u32 %p_jw, 1, 0; into the array. The driver's own dump shows it (/tmp/failed-ptx-sm_121-_visible__entry_iq3_s_gemv_warp_reduce_.ptx, lines 6–8):

.global .align 4 .u32 iq3s_grid_g[512] = {
    .reg .pred %p_jw;
    setp.ne.u32 %p_jw, 1, 0;

Scope

Devices. The pass runs when major >= 12 (driver/module.rs:206 and :293): sm_120 and sm_121. sm_89 (lambda, RTX 4090) never runs it, so an sm_89 A/B cannot see this defect.
Kernels. Any kernel with a = { global table ahead of its .entry that also has at least one backward branch. The GPU whitelist admits a quant TYPE on one model's shapes — admission must key on per-shape device conformance (census-derived, #3945) #3968 harness on GB10 measured IQ3_S and IQ2_S failing; IQ2_XXS, IQ3_XXS, IQ4_NL and IQ4_XS loading.
CERTIFIED 0.69.1 models are NOT affected. Tensor types from apr tensors --json:
- Qwen3.5-2B-Q4_K_M: F32, Q4_K, Q5_K, Q6_K, Q8_0;
- 4B-Q4_K_M: the same set;
- 4B-UD-Q4_K_XL: F32, F16, Q4_K, Q5_K, Q6_K, Q8_0, IQ4_XS (IQ4_XS loads).
- End to end, the X2 CRUX smoke was 134/134 GREEN on gx10.
Why no test caught it. Every test in ptx_patch.rs begins its module with .visible .entry. None puts a global table before the entry, and the rewriter's tests run only in the gx10 nightly filter (aprender-gpu: driver, kernels, memory AND ptx are all #[cfg(feature = "cuda")] — 2,160+ tests incl. the GH-480 rewriter's and the PTX emitter's run ONLY in the gx10 nightly's perf053 filter #3067).

Fix and must-RED case

Fix direction (dd): set in_body only for the { that opens a .entry/.func body. Track having seen .entry/.func, or never treat a line containing = as a body opener.
Must-RED case for ptx_patch's own tests: a module with .global .u32 t[2] = {\n 1, 2\n}; BEFORE an .entry that has a backward branch. Today the .reg lands inside the initializer; after the fix it must land after the entry's {.
Re-verify on gx10: bash scripts/gpu_shape_conformance.sh gx10 (feat/3968-gpu-whitelist-shape-conformance), rows 21/22.

Refs #3968 #3067 #4004

Answering the open question (aprender-dd): the IQ3_S and IQ2_S admissions were measured ONLY on sm_89, where patch_backward_branches_sm121 never runs, so no admission measurement could see this defect. From the admitting tests' own doc comments in crates/aprender-serve/src/cuda/gemv_entry_name_tests_3477.rs:

IQ3_S (apr serve returns HTTP 500 on every route for an fp16 .apr model that run/chat/golden all handle correctly #3884): "EXACT against iq_parallel_matvec on the same bytes, RTX 4090 sm_89".
IQ2_S (IQ2_S (ggml 22) has no CUDA GEMV — 0.69.1 release blocker #3953, with IQ3_XXS and Q2_K): "Bar 1e-5, NaN-prefilled, RTX 4090 sm_89, card exclusive".

That is the generator #3968 exists for: admission measured on one architecture licensed a type on every architecture, including the one (cc ≥ 12) whose driver path rewrites the PTX. The census-driven harness runs per host (scripts/gpu_shape_conformance.sh <host>), and its guard requires a passing row on EVERY host that holds the shape.

docs(release): 0.69.1 known issue — IQ3_S/IQ2_S CUDA GEMV do not load…

evidence(cuda): gx10 shape-conformance receipt @c603a2405 — 160/163 g…

evidence(cuda): lambda shape-conformance receipt @65f752cd5 — 163/163…

feat(cuda): IQ3_S/IQ2_S are refused on cc >= 12 until #4096 — the whi…

Milestoned 0.70.0 (cop): not a 0.69.1 release blocker (fails closed; no certified model affected). The fix is ptx_patch.rs in_body detection plus a must-RED on the iq3s_grid_g global initializer. Owner after publish: aprender-f5.

evidence(cuda): gx10 receipt @5c1df1b3e — cc 12, IQ3_S/IQ2_S excluded…

fix(#3968): quorum lane 1 — apr capability renders the cc>=12 exclusi…

Claimed by aprender-cb (assigned by the cop). I'll commit onto batch/0.70.0. Scope: the ptx_patch.rs in_body detection, a must-RED test on the iq3s_grid_g initializer, and lifting the #3968 cc>=12 exclusion with a gx10 receipt.

fix(cuda): GH-480 patch pass declared its predicate inside a .global …

Fix landed on batch/0.70.0 @ f049cf9 (fast-forward; the PR cap is full, so it rides the batch).

Root cause. patch_backward_branches_sm121 (GH-480, crates/aprender-gpu/src/driver/ptx_patch.rs) runs only when cc>=12, so sm_89 never exercised it. It took the first line ending in { as the kernel body. In IQ3_S/IQ2_S that line is the .global .u32 iq3s_grid_g[512] = { / iq2s_grid_g[2048] codebook initializer, so .reg .pred %p_jw; setp… landed inside the array and ptxas failed with Parsing error near '.reg'.

Fix. Brace-depth scope scan: the decl goes at depth 1 of each .entry/.func body, once per function. Labels are keyed per function, and a prototype ; disarms the header. The two latent bugs are fixed too (one decl per module; one global label map).

Evidence (measured):

Tests: 4 new unit tests in driver::ptx_patch, 12/12 pass with --features cuda. Against the old pass body, 3 of them fail (mut rc=101: initializer, per-function, prototype).
Real kernels: dumped all 13 layout.rs GEMV generators from rc/0.69.3-rc.2 at k=n=4096. Patched them with the old and the new pass, then ran ptxas -arch=sm_121 on gx10 (CUDA 13.3, cuda_13.3.r13.3/compiler.38244171_0):
- raw (unpatched): 13/13 OK
- old pass: iq2_s and iq3_s fail with line 8/9; fatal: Parsing error near '.reg'
- new pass: 13/13 OK. Every file has decl=1, and 1–2 @%p_jw bra sites were patched, so the pass did engage.
Pre-push: fmt ok, aprender-contracts --lib 1701 pass, cargo deny advisories ok.

Not done / out of scope:

The GPU whitelist admits a quant TYPE on one model's shapes — admission must key on per-shape device conformance (census-derived, #3945) #3968 cc>=12 IQ3_S/IQ2_S exclusion exists only on the unmerged feat/3968-gpu-whitelist-shape-conformance, so there was nothing to lift on batch/0.70.0.
An end-to-end apr run on gx10 was not run: gx10 root is 98% full (23G free, below the 50G watchdog floor), so I did not build apr there. The ptxas-on-sm_121 result is a module-load proof, not an inference-correctness proof.

#4096 e2e on gx10 (GB10, sm_121): apr run --gpu --json, greedy, 32 tokens, 3 prompts. The fix is f049cf9 on batch/0.70.0, cherry-picked on gx10 as 84be0ac70.

model	binary	runs on GPU (`backend.ran=gpu`, `fell_back=false`)	result
Qwen2.5-0.5B IQ3_M	RED `15c3032`	0/3	rc=14 "gpu was forced… ran on CPU"; ptxas `line 8; Parsing error near '.reg'`
Qwen2.5-0.5B IQ3_M	GREEN 84be0ac70	3/3	no ptxas error
Qwen2.5-0.5B IQ2_M	RED	0/3	same ptxas error
Qwen2.5-0.5B IQ2_M	GREEN	3/3	tokens identical to sm_89 (RTX 4090) on all 3 prompts

The IQ3_S/IQ2_S kernels now load and run on sm_121. Before the fix they could not be JIT'd at all.

Two findings outside this issue:

sm_121 GPU and CPU disagree even without IQ types. Test model: qwen2.5-coder-0.5b Q4_K_M on GB10. Matching prefix between the GPU and CPU runs: p1 1/1, p2 0/32, p3 32/32. IQ3_M on GB10 also diverges from CPU at tokens 6 and 5 on p2 and p3, while sm_89 matches CPU 32/32. Since plain Q4_K (IQ3_M ffn_down in blk 0/1/10 is Q4_K) shows the same pattern, this looks like sm_121/Q4_K numerics, not the backward-branch patch.
The real-model device A/B harness fails on both archs. real_model_iq_ab leaves "34 of 4864 rows never written" (IQ3_S, blk.11.ffn_down) and "35 of 4864" (IQ2_S, blk.3.ffn_down), identically on sm_89 and sm_121. So the failure is pre-existing and not arch-dependent. The likely cause is the harness's dims (k=896 is not a multiple of 256). The synthetic device A/B passes 14/14 on both archs (iq3_s, iq2_s, iq4_nl, q5_1, q2_k, bf16, iq2_xxs, iq3_xxs).

Neither of these blocks #4096. Both need their own tickets; I'm leaving that call to the cop.

Follow-up: the Q4_K_M control on sm_89 (RTX 4090, lambda), same model/prompts, GPU vs CPU prefix. p1 1/1, p2 32/32, p3 7/26 (sm_89 stops early at EOS). On sm_121 the same prompts give p1 1/1, p2 0/32, p3 32/32. So both GPU archs drift from CPU on plain Q4_K_M, but only sm_121 splits at the first token (p2). That's still one prompt, so it is not yet a confirmed sm_121 defect; it needs more prompts and a logit-margin check before a ticket. It remains out of scope for #4096.

Status (la-0.71, owner per label), 2026-10-12.

Fix is on main. The ptx_patch.rs scope detection and its must-RED tests (global_initializer_is_not_a_function_body_4096, one_line_initializer_before_entry_is_module_scope, every_function_declares_its_own_predicate_and_labels_are_function_scoped, a_prototype_does_not_arm_the_next_brace) arrived with 989cb01 (fold(0.70): #4587+#4588 PMAT-4502, #4620 #3715, #4614 #4616, #4640 PMAT-937 (C221 fat_driver fix) #4655). b2-gpu.yml runs them in cargo test -p aprender-gpu --lib --features cuda.
Device proof on file. The 2026-09-24 gx10 run above (GB10, sm_121): IQ3_M and IQ2_M went from 0/3 to 3/3 on the GPU, with no ptxas error, and IQ2_M tokens identical to sm_89.
Fresh proof on a current binary: pending. gx10's installed apr 0.71.1 (e89fb30c4) is on main and contains the fix (git merge-base --is-ancestor 989cb012e5 e89fb30c4 = 0). But gx10's root disk reads 90%, so under the no-load-at-≥90% rule I have not run on it. Once it is below 90%, the rerun is the same three-prompt apr run --gpu --json on Qwen2.5-0.5B-Instruct-IQ3_M.gguf and cb-4096-Qwen2.5-0.5B-IQ2_M.gguf (both already on gx10), pass = backend.ran=gpu, fell_back=false, 3/3.
No code change is needed and there is no branch. Closure is the cop's call.

Agent: la-0.71

P0 to P1 (2026-10-14 P0 triage): P0 now means it blocks the next release or carries the operator's own words on it; this issue does neither, so it stays on the 0.71.0 milestone at P1.

Re-parent: former parent #3953 is closed. Moved under open epic #3994 (same 0.71 train).

merged into #5312 (master ticket for epic #3994; triage rule 7b, 2026-10-14). It is a checklist row there; reopen this issue when work on it starts or its PR opens.

Moved: this sub-ticket is now a row in #5323 (CUDA quantized kernels: correct block layouts, cache keys and precision), split out of #5312 by theme.

Navigation Menu

Repository navigation

P0: IQ3_S/IQ2_S GEMV kernels fail to load on cc>=12 (GB10/Blackwell) — GH-480 PTX patch injects .reg inside a global array initializer #4096

P0: IQ3_S / IQ2_S GEMV kernels cannot load on compute capability ≥ 12 (GB10 sm_121, RTX 50xx sm_120)

Reproduced on the 0.69.1 release binary (X2)

Root cause (`crates/aprender-gpu/src/driver/ptx_patch.rs`, at `d8a6df5`)

Scope

Fix and must-RED case

15 remaining items

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

P0: IQ3_S/IQ2_S GEMV kernels fail to load on cc>=12 (GB10/Blackwell) — GH-480 PTX patch injects .reg inside a global array initializer #4096

Description

P0: IQ3_S / IQ2_S GEMV kernels cannot load on compute capability ≥ 12 (GB10 sm_121, RTX 50xx sm_120)

Reproduced on the 0.69.1 release binary (X2)

Root cause (crates/aprender-gpu/src/driver/ptx_patch.rs, at d8a6df5)

Scope

Fix and must-RED case

Activity

noahgift commented on Sep 23, 2026

noahgift commented on Sep 23, 2026

noahgift commented on Sep 24, 2026

noahgift commented on Sep 24, 2026

noahgift commented on Sep 24, 2026

noahgift commented on Sep 24, 2026

15 remaining items

noahgift commented on Oct 12, 2026

noahgift commented on Oct 14, 2026

noahgift commented on Oct 14, 2026

noahgift commented on Oct 14, 2026

noahgift commented on Oct 14, 2026

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions

Root cause (`crates/aprender-gpu/src/driver/ptx_patch.rs`, at `d8a6df5`)