fix(intrinsics): admit FP16 m16n8k8 MMA on sm_75

[praveenperera]

What this adds

mma_m16n8k8_f16_f16 and mma_m16n8k8_f32_f16 now carry the PTX ISA floor for .f16 m16n8k8, PTX 6.5 and sm_75, on both backends, so a Turing kernel can call them as intrinsics instead of going through ptx_asm!. The BF16 and TF32 m16n8k8 forms and the m16n8k16 form keep PTX 7.0 and sm_80.

before: cargo oxide build --arch sm_75
        error: CUDA target sm_75 cannot lower generated intrinsic `mma_m16n8k8_f16_f16` (`v1:i0521`); requires sm_80 or newer

after:  cargo oxide build --arch sm_75
        .version 6.5 / .target sm_75
        mma.sync.aligned.m16n8k8.row.col.f16.f16.f16.f16 {...}
        mma.sync.aligned.m16n8k8.row.col.f32.f16.f16.f32 {...}
        ptxas -arch=sm_75: HMMA.1688.F16, HMMA.1688.F32

#[kernel]
fn turing_mma(out: &mut [f32], a: &[u32], b: &[u32], c: &[f32]) {
    let lane = thread::index_1d().get();
    let a_frag = [a[lane * 2], a[lane * 2 + 1]];
    let c_frag = [c[lane * 4], c[lane * 4 + 1], c[lane * 4 + 2], c[lane * 4 + 3]];
    // SAFETY: one full warp; every lane executes the same instruction.
    let d = unsafe { wmma::mma_m16n8k8_f32_f16(c_frag, a_frag, b[lane]) };
    out[lane * 4] = d[0];
}

What we fixed

• The two FP16 m16n8k8 recipes carried the admission-wide PTX 7.0 / sm_80 floor. The PTX ISA introduces .f16 m16n8k8 at PTX 6.5 for sm_75; .bf16 and .tf32 m16n8k8 and .f16 m16n8k16 arrive at PTX 7.0 for sm_80.
• The floor moved at its one source, the recipe table. The admission copies it into the policy and both backend lowerings, so the catalog, the generated target requirements, the device stub docs and the reference table follow by regeneration.
• LLVM's selection predicates admit the BF16 and TF32 m16n8k8 forms at sm_75 as well, and ptxas rejects them there (Feature '.bf16' requires .target sm_80 or higher). The floor follows the ISA per form, not the shared predicate, and the generator test asserts that a BF16 policy lowered to sm_75 is refused.
• Landing work on top of the submitted change: the Turing LLVM evidence is recorded under the current rust-llvm-23.1.3-aab11ee0 profile from real runs of the nightly-2026-10-11 sysroot llc, ptxas 13.3 V13.3.73 and the sysroot llvm-as, llvm-dis and opt (the sm_75 PTX is byte-identical to the submitted record; the libNVVM evidence is the author's, unchanged); the generated summary no longer calls the two Turing forms Ampere forms; both recipes state the ISA rule; the codegen negative test covers every sm_80 member of the admission at sm_75; main is merged and the catalog outputs regenerated with the book stamp refreshed.

Floors after this change

Intrinsic	PTX	SM
mma_m16n8k8_f16_f16, mma_m16n8k8_f32_f16	6.5	sm_75
mma_m16n8k8_f32_bf16	7.0	sm_80
mma_m16n8k4_f32_tf32, mma_m16n8k8_f32_tf32	7.0	sm_80
mma_m16n8k16_f16_f16	7.0	sm_80

Implementation

• Recipe table: minimum_ptx: "6.5", minimum_sm: "sm_75" on the two FP16 m16n8k8 recipes, with the ISA rule in a comment on both; the other admission members are unchanged.
• Evidence: both register-mma-ampere-float files gain a turing_dense_f16_mma_m16n8k8 fixture with per-variant facts (sm_75 / +ptx65 for LLVM, compute_75 for libNVVM), appended after the Ampere fixture so the refresh ledger's pointers stay valid. ampere_dense_float_mma_all keeps its sm_80 stages and covers the remaining three forms (coverage_count 5 to 3). The generator requires the selected backend_codegen stage to sit at the catalog floor, which is why the fixture exists.
• LLVM stages: declaration canonicalization, llc -mcpu=sm_75 -mattr=+ptx65 on a two-form kernel fixture, ptxas -arch=sm_75. libNVVM stages: compute_75 PTX and LTOIR, ptxas, nvJitLink -arch=sm_75 -lto. runtime_validation stays unexecuted.
• Generated summary: Executes the <id> dense floating-point MMA form. for the five records of the admission, so the stub docs of the Turing forms no longer read "Ampere" above "Requires PTX 6.5 and sm_75+".
• Tests: the generator tests check a per-form floor table and fail closed (the FP16 forms at the old floor are rejected; a BF16 policy at sm_75 is rejected). The codegen test generated_m16n8k8_mma_floors_admit_turing_only_for_fp16 covers both backends: PTX 6.5 requirement, sm_72 rejected, sm_75 and sm_80 accepted, and the four sm_80 members (i0522, i0107, i0520, i0524) refused at sm_75 with an error naming the intrinsic and sm_80 or newer.
• Regenerated outputs: catalog.json (two records change floors and evidence strings, three change summary only), the generated target tables (7.0 to 6.5 and sm_80 to sm_75 for i0521 and i0523 only), stub docs, reference table, catalog SHA banners across the generated crates, and the book stamp. abi-v1.toml, upstream.lock and the overlay shards are untouched.

Verification

Checked at 9fe89abb4 (nightly-2026-10-11, RTX 5090, CUDA 13.4; ptxas 13.3 for the evidence):

• ✅ cargo run -p cuda-intrinsics-gen -- check; probe --all --skip-terminal --per-target with both FP16 forms selecting sm_75 +ptx65 and BF16 sm_80 +ptx70; check-abi-history with every base entry preserved and the total unchanged; a second generate is a no-op; the book catalog stamp guard passes.
• ✅ cargo test -p cuda-intrinsics-gen --all-targets, cargo test -p cuda-oxide-codegen --lib, clippy with warnings as errors on both, rustdoc, cargo oxide fmt --check, the 18 repository guards, the example lock check, the book with warnings as errors.
• ✅ Evidence reproduced: every recorded size and SHA-256 of the Turing fixture reproduces from its inputs with the nightly-2026-10-11 sysroot llc (LLVM 23.1.3), ptxas 13.3 V13.3.73 and the sysroot llvm-as, llvm-dis and opt; the sm_75 PTX is byte-identical to the submitted record.
• ✅ Fail-before: with the two recipe floors reverted, three generator tests and check fail; main's generated target table fails the codegen test at the i0521 assertion.
• ✅ RTX 5090, JIT from the sm_75 PTX: a one-warp kernel built with --arch sm_75 matches a host reference (f32 accumulate bit-exact, f16 accumulate within 1.7e-3 on random inputs) with memcheck and synccheck clean; ptxas -arch=sm_75 turns it into HMMA.1688.F16 and HMMA.1688.F32; generated_intrinsics on sm_120a passes as a control. At the submitted head the same kernel also passed through the libNVVM route.
• ✅ Negatives: mma_m16n8k8_f32_bf16 at sm_75 is refused by the compiler ("requires sm_80 or newer") and by ptxas ("Feature '.bf16' requires .target sm_80 or higher"); sm_70 and sm_72 are refused naming the intrinsic.
• ✅ Tesla T4 (author's run, driver 610.57.04, CUDA 13.3): the kernel built with --arch sm_75 ran both forms in all three device-artifact modes; every output matched a host reference, bit-exact for small integers and within 6.5e-4 (f16 accumulate) and 1.5e-7 (f32 accumulate) for random inputs in [-1, 1]; SASS HMMA.1688.F16 and HMMA.1688.F32.

Caveats

• The maintainer's box has no Turing GPU; the T4 numbers are the author's.
• The floor below sm_75 is enforced by the compiler only: CUDA 13 ptxas defines no sm_70 or sm_72 target, so no assembler rejection can be recorded for them.
• Every sm_75 run used an explicit --arch sm_75. Without it, target selection accepts the detected GPU once the generated requirement is satisfied; that path is covered by code reading only.
• The T4 execution cannot enter the evidence files: the admission requires runtime_validation to stay unexecuted (intrinsics-gen: runtime_validation = "executed" is unreachable for compact admissions (30 guards, no Runtime-stage producer) #1298).
• mma_m16n8k8_f32_tf32 at sm_75 is refused by the module-level feature detector ("cannot lower detected feature Sm80") before the per-intrinsic floor is consulted, so its message differs from the BF16 case.
• Follow-ups outside this PR: the hand-written mnemonic scanner in cuda-oxide-codegen has no entry for the two FP16 m16n8k8 forms, so external LLVM IR fed through the pipeline at --arch sm_75 still gets the recorded PTX 6.3 floor (cuda-oxide kernels are unaffected: the generated requirement raises the module to 6.5); and no in-tree example builds at sm_75 yet.

Closes #1467.

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[nihalpasham]

/ok to test ed558e1

[nihalpasham]

/ok to test 9fe89ab

[nihalpasham]

Thanks, this is the right fix: the floor moves at its one source, the recipe table, and the per-form split matches the PTX ISA. I merged main in and layered the landing work on top:

• Re-recorded the LLVM evidence for the Turing fixture under the current rust-llvm-23.1.3-aab11ee0 profile with the nightly-2026-10-11 sysroot llc and ptxas 13.3 V13.3.73. The sm_75 PTX is byte for byte what you recorded, so only the tool provenance and the cubin hash changed. Your libNVVM evidence is unchanged.
• Reworded the generated summary so the two sm_75 forms are not called Ampere forms right above a sm_75+ requirement, and put the ISA rule on both recipes.
• Added the m16n8k4 TF32 and m16n8k16 F16 forms to the codegen negative test, so every sm_80 member of the admission is checked at sm_75.
• Regenerated the catalog outputs and refreshed the book stamp.

At 9fe89abb4, the generator gates, both test suites, clippy, fmt, the guards, the book and all CI checks passed. On the 5090 a one-warp kernel built with --arch sm_75 matched a host reference with memcheck and synccheck clean, and ptxas -arch=sm_75 turns it into HMMA.1688.F16 and HMMA.1688.F32. No Turing GPU here, so your T4 run is the hardware evidence.

Closes #1467. Below sm_75 the compiler alone enforces the floor: CUDA 13 ptxas has no sm_70 or sm_72 target to record a rejection against.