bug(cuda): 12 committed PTX files regressed to .version 9.1 — unloadable on pre-13.1 drivers (#124 regression, no guard)

[jamesburton]

Observation

12 of the 57 committed PTX files are at .version 9.1, silently reverting the portability fix from #124. The other 45 are correctly at .version 8.7.

attention.ptx  attention_f32.ptx  attention_flash_mma_decode_gqa_split.ptx
conv1d_causal.ptx  dequant.ptx  dequant_bf16_mxfp4.ptx  dequant_iq1.ptx
dequant_iq3.ptx  dequant_iquants.ptx  dequant_pq2_0.ptx
gated_delta_net_scan.ptx  pq2_0_gemv.ptx

.version 9.1 PTX fails to load with CUDA_ERROR_UNSUPPORTED_PTX_VERSION on any driver older than CUDA 13.1 — the exact failure #124 was raised to fix, after a Kaggle T4 session on CUDA 13.0 could not load our PTX.

.target sm_75 is not part of this — that is the intended arch policy and every file agrees on it (with the documented sm_86 overrides for the mma.sync kernels). Only the ISA version drifted.

How it happened

Anyone rebuilding a kernel with the CUDA 13.1 toolkit re-raises the ISA version for that file only, so the regression arrives one file at a time and is invisible in review — the diff is a wall of regenerated PTX in which line 9 changing is not conspicuous. Bisect of dequant.ptx:

commit	.version
db7b1f3a (#124, rebuilt at 12.8)	8.7
5d724b8d (#254 Q4_0/Q4_1 fix)	9.1 ← regressed here
a9b3700e (#311 Q3_K fix)	9.1 (inherited)

There is no guard on dev. A PtxTargetTests exists, but only on the issue/156-q2k-cpu-q3k-q41-q51-dequant branch (open upstream PR #161), and it asserts the target arch, not the ISA version — so it would not have caught this even once merged.

The rebuild is possible — a prior conclusion to the contrary was wrong

An investigation concluded regeneration was impossible on the T5500, because C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8 is a stub and CUDA 13.1 rejects compute_61.

The real CUDA 12.8 toolkit is at E:\CUDA_v12.8.1\bin\nvcc.exe. where nvcc on that box returns both; the check stopped at the first hit. cl.exe is also not on the default PATH but MSVC is present — a VS dev shell is needed first. This is the same toolkit #124 used.

Acceptance criteria

• All 12 files rebuilt with the CUDA 12.8 toolkit so every committed PTX is .version 8.7.
• Arch targets unchanged (sm_75, with the existing sm_86 overrides preserved).
• Entry-point sets verified identical before/after — a rebuild must not silently add, drop or rename a kernel.
• A guard test asserting .version <= 8.7 across every file in native/ptx/, so the next 13.1 rebuild fails CI instead of shipping. This is the part that stops it recurring; the arch-only check on PQ2_0 decode: eliminate F32<->F16 conversion launches + ALU-bound ternary unpack reduction #161's branch is not sufficient.
• native/build_ptx documents which toolkit to use and why, including the E: path.

Note on scope

dequant.ptx currently contains the pre-fix Q3_K kernel on the #161 branch (see that PR), and the fixed one on dev. Whoever does this rebuild should confirm which source they are compiling from, and that the resulting Q3_K path matches the corrected .cu.

References

• cuda(native): rebuild committed PTX at CUDA 12.8 to fix UNSUPPORTED_PTX_VERSION on pre-13.1 drivers #124 — the original portability fix and its rationale
• bug(vulkan): Bielik-1.5B Q3_K_M fails CPU parity at prefill step 0 (L-inf 5.09 vs 3.00) while Q4_K_M passes #311 — Q3_K correctness fix; its CUDA half is unverifiable at runtime until PTX is trustworthy
• upstream cuda(quant): Q2_K CPU dequant + Q3_K / Q4_1 / Q5_1 CPU+CUDA dequant kkokosa/dotLLM#161 — carries PtxTargetTests (arch only)