Skip to content

feat(intrinsics): admit bf16x2 to FP8x2 pack conversions - #1354

Merged
nihalpasham merged 9 commits into
NVIDIA:mainfrom
YeonwooSung:feat/bf16x2-to-fp8x2-pack
Oct 19, 2026
Merged

nihalpasham merged 9 commits into
NVIDIA:mainfrom
YeonwooSung:feat/bf16x2-to-fp8x2-pack

Conversation

@YeonwooSung

@YeonwooSung YeonwooSung commented Sep 28, 2026 •

Copy link
Copy Markdown
Contributor

What this adds

Four Blackwell-family pack conversions from packed bf16x2 to packed FP8, admitted to the intrinsic catalog as inline-PTX records: cvt_rn_satfinite_e4m3x2_bf16x2, cvt_rn_satfinite_relu_e4m3x2_bf16x2, cvt_rn_satfinite_e5m2x2_bf16x2 and cvt_rn_satfinite_relu_e5m2x2_bf16x2. The generator emits the cuda_device::convert stubs, dialect-nvvm ops, importer arms, mir-lower lowerings, ABI ledger entries i1030 to i1033 and the reference rows, the same way as the f16x2 siblings. cuda-target-spec gains the +ptx91 feature these instructions need.

before: a kernel holding bf16 pairs unpacks to two f32 and packs those
  cvt.f32.bf16 + cvt.f32.bf16 + cvt.rn.satfinite.e4m3x2.f32     3 instructions, 2 extra registers per pair
after:
  cvt.rn.satfinite.e4m3x2.bf16x2 %rs1, %r1;                       1 instruction, PTX 9.1, Blackwell families
use cuda_device::convert::{cvt_rn_satfinite_e4m3x2_bf16x2, cvt_rn_satfinite_relu_e5m2x2_bf16x2};

// packed bf16x2 bits in, packed FP8x2 bits out; the low bf16 half lands in the low byte
let e4m3: u16 = cvt_rn_satfinite_e4m3x2_bf16x2(packed_bf16x2);
let e5m2_relu: u16 = cvt_rn_satfinite_relu_e5m2x2_bf16x2(packed_bf16x2);

Build with cargo oxide run <example> --arch sm_120a (or sm_120f). CUDA 13.4's own cuda_fp8.hpp emits the same instruction for __nv_cvt_bfloat16raw2_to_fp8x2 under family-specific compilation.

What we fixed

  • Recorded the LLVM evidence under the current rust-llvm-23.1.3-aab11ee0 profile with the typed-NVVM stages and the ptxas assembly the packed-conversion family requires, and removed the bf16x2 exemption from the evidence validator. The pinned llc selects llvm.nvvm.bf16x2.to.* at sm_100f and sm_120a with +ptx91.
  • Replaced the libNVVM evidence with a real libNVVM 13.3.73 compile, ptxas assembly and nvJitLink link; the submitted file reused the f16x2 sibling's header around an llc run.
  • Widened the target set from sm_100f|sm_110f|sm_120f to the family list the catalog uses for Blackwell records. The compiler matches family targets on the literal f suffix, so the short list refused --arch sm_120a and sent an arch-less cargo oxide run on an RTX 5090 to sm_100f, which failed at module load.
  • Added the PTX 9.1 needs LLVM 23 rule to the PTX ISA gate. LLVM 22 knows +ptx90 only; with CUDA_OXIDE_LLC=llc-22 the build was silent, the PTX said .version 8.8, and the JIT failed at load. The examples-compile lane's LLVM floor moves from 22 to 23 for the same reason.
  • Merged main and regenerated, refreshed the book catalog stamp and counts, added the negative admission cases the f16x2 test has, and extended generated_intrinsics_blackwell with the four bf16x2 forms.

Target contract

  • PTX 9.1 floor. Accepted targets: sm_100a, sm_100f, sm_103a, sm_103f, sm_110a, sm_110f, sm_120a, sm_120f, sm_121a, sm_121f. The PTX ISA admits family features on f-targets and on a-targets of the same family; ptxas 13.4 and the pinned llc accept the same set.
  • Refused at compile time: unsuffixed sm_100, sm_103, sm_110, sm_120 and sm_121, sm_90a and older, and any llc older than LLVM 23.
  • cargo oxide run with no --arch selects the detected device (sm_120a on a 5090); cargo oxide build with no --arch compiles for the first catalog alternative, as for every family-gated record, so pass --arch for a specific device.
Implementation
  • intrinsics/overlay/packed_conversion_fp8_bf16x2.toml (schema 65) is a compact shard: fp8_formats = ["e4m3x2", "e5m2x2"], relu_variants = true, product_count = 4, runtime_validation = "unexecuted"; listed in overlay.toml.
  • PackedConversionFp8Bf16x2Admission fails closed on any axis deviation and builds records through the shared packed_conversion_overlay_record; four closed recipes carry the ids, ABI ids, operation keys, the v2bf16 -> i16 typed signature and the mnemonics.
  • PackedConversionSourceFormat::Bf16x2 extends every render and resolve match beside F16x2: 32-bit u32 source, u16 result, =h,r constraint, identity half order; the generated mir-lower impl is a convert_generated_packed_unary call like its siblings.
  • packed_conversion_floor returns the 9.1 floor with no minimum SM and the family target list through a named const with the derivation in its comment.
  • cuda-target-spec adds 91 to the PTX feature vocabulary; the gate in cuda-oxide-codegen is a table of PTX floor to LLVM major (9.0 needs 22, 9.1 needs 23).
  • Evidence: rust-llvm-23.1.3-aab11ee0-packed-conversion-fp8-bf16x2.json and cuda-13.3-libnvvm-13.3.73-packed-conversion-fp8-bf16x2.json, recorded from real tool runs.

Verification

Checked at 89e24e263 (nightly-2026-10-11, RTX 5090):

  • ✅ cuda-intrinsics-gen check current, probe --all --skip-terminal --per-target green, check-abi-history with the 1029 base entries preserved and 1033 total; 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 (the gate test fails without the 9.1 rule), the touched crates' clippy with warnings as errors, rustdoc, fmt, the 18 guards and the lock check.
  • ✅ cargo oxide run generated_intrinsics_blackwell --arch sm_120a: all 16 bf16x2 checks match (0x4038/0x403c, 0x7eb8/0x7e00, +Inf and NaN cases), memcheck 0 errors; the PTX is .version 9.1 with one cvt.rn.satfinite[.relu].{e4m3x2,e5m2x2}.bf16x2 per form.
  • ✅ Arch-less cargo oxide run on the 5090 selects sm_120a and passes; CUDA_OXIDE_LLC=llc-22 is refused at compile time with the LLVM 23 message.
  • ✅ Review runs at the submitted head: all four forms over every bf16 pattern in each lane, 0 mismatches against CUDA 13.4's cuda_fp8.h host conversion, through the llc and libNVVM routes; ptxas 13.4 accepts the instruction on the ten family targets and rejects unsuffixed sm_100 and sm_120.

Caveats

  • runtime_validation stays unexecuted in the catalog: compact admissions have no runtime stage to record (intrinsics-gen: runtime_validation = "executed" is unreachable for compact admissions (30 guards, no Runtime-stage producer) #1298), so the GPU results live in this description and in the Blackwell example.
  • Pack only. The pinned LLVM import has no unscaled unpack record (cvt.rn.bf16x2.e4m3x2 is PTX 9.2).
  • Family-gated: there is no sm_89 or sm_90 form of this instruction; the f32-pair conversions remain the portable route.
  • probe --candidate cannot record evidence for this family (a pre-existing panic in the packed renderer); the evidence was recorded with the same process as the f16x2 refresh.

Admit the four selectionless bf16x2 pack records as inline PTX. They require
PTX ISA 9.1 on sm_100f, sm_110f, and sm_120f. Runtime validation stays
unexecuted because ptxas and libNVVM are not available here.

Signed-off-by: YeonwooSung <neos960518@gmail.com>
The SIMT tree now lives under cuda-oxide/. Directory-rename detection
moved the bf16x2-to-FP8x2 evidence, overlay shard, and probes with
intrinsics/. The four selectionless pack records remain the only new
admission. Regenerating the catalog was a no-op.

Signed-off-by: YeonwooSung <neos960518@gmail.com>
@copy-pr-bot

copy-pr-bot Bot commented Oct 5, 2026

Copy link
Copy Markdown

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 nihalpasham added the cuda-oxide SIMT programming model: rustc backend, cargo oxide, cuda-device, cuda-host, book label Oct 11, 2026
Take main's copy of every generated, probe, evidence and lock file that conflicted. The generated outputs are regenerated in a later commit once the bf16x2 evidence is re-recorded under the current LLVM profile. The overlay shard, ledger entries, generator and target-spec changes keep the branch's side.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
The PTX vocabulary now contains +ptx91, so a 9.1 floor can reach an LLVM 22 llc. LLVM 22 ignores the feature with a warning and emits PTX 8.8; the build passed and the module failed to load. The gate refuses PTX 9.1 or newer on anything older than LLVM 23. The fixture for a floor newer than cuda-oxide can request moves to 9.2, and the floors test covers the 9.1 rule.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
…ofiles

The four bf16x2 to packed-FP8 records now meet the packed-conversion family contract instead of being exempted from it. The LLVM evidence is re-recorded with the pinned nightly-2026-10-11 llc and ptxas 13.3.73 at sm_120f with +ptx91: typed declaration canonicalization, typed and inline codegen, and ptxas assembly of both routes, which produce the same cubin. The libNVVM evidence is a real libNVVM 13.3.73 compile at compute_120f, a ptxas assembly and an nvJitLink 13.3 link of the LTOIR; the previous file reused another record's header around an llc run.

The target list is the family-complete Blackwell set the catalog uses for its other family records, because cuda-oxide-codegen matches a family target only on the f suffix and the short list refused sm_120a. The admission test gains the fail-closed cases, the f16x2 field gets its serde default back, the generated outputs are regenerated, and the book catalog stamp moves with them.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
…ackwell

Adds a kernel that calls the four bf16x2 to packed-FP8 conversions and four host cases checked bit for bit, including +inf saturation and NaN. Verified on an RTX 5090 at sm_120a; compile-only in CI as before.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
generated_intrinsics_blackwell now calls the four bf16x2 to packed-FP8
conversions, which are PTX 9.1 forms, so its PTX header moved from
.version 8.7 to .version 9.1 and the blackwell-compile lane's header
pin went red while every other assertion still held. Pin the new floor
and prove the new instructions the way the lane proves the f32 packs:
four distinct cvt.rn.satfinite[.relu].{e4m3x2,e5m2x2}.bf16x2
instructions in the llc-route PTX and four "=h,r" inline-asm calls in
the libNVVM-route NVVM IR.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
generated_intrinsics_blackwell now needs +ptx91 for the bf16x2 to
packed-FP8 conversions. llc-22 knows +ptx90 at most, and the codegen
gate refuses a PTX 9.1 request below LLVM 23 at compile time, so the
compile-only lane cannot build that example under its LLVM 22 pin.
Install LLVM 23 from apt.llvm.org instead (the llvm-toolchain-*-23
suites serve 23.1.4 for noble and resolute) and point both
CUDA_OXIDE_LLC sites at llc-23, following the previous floor bump from
21 to 22. The floor comment names the new reason and keeps the old one.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
The inline-route ptx_assembly detail of the two e5m2x2 records repeated
the e4m3x2 records' negative control. Retargeting each form's inline
PTX to sm_120 with ptxas 13.3.73 rejects the e5m2x2 forms with
"Feature 'cvt.e5m2x2.bf16x2' not supported on .target 'sm_120'", and
the two details now quote that message. The catalog carries the detail
text, so the generated outputs and the book's catalog stamp move with
it; no route, target or signature changes.

Signed-off-by: nihalpasham <nihalp@nvidia.com>
@nihalpasham nihalpasham added enhancement New feature or request intrinsics Device intrinsics and libdevice math mappings cuda-feature A CUDA hardware/toolkit feature (TMA, tcgen05, WGMMA) device-apis User-facing device-side and kernel-authoring APIs (cuda-device, cuda-macros) build-related Build scripts, toolchain discovery, or compilation setup labels Oct 19, 2026
@nihalpasham

Copy link
Copy Markdown
Collaborator

/ok to test 89e24e2

@nihalpasham

Copy link
Copy Markdown
Collaborator

Thanks, the four records are right and the generated code is correct: on the 5090 all four forms are bit for bit equal to CUDA's own cuda_fp8.h conversion. The gaps were in the evidence and the gates, so I layered those on top:

  • Re-recorded the LLVM evidence under the current toolchain profile with the typed stages the family requires, and replaced the libNVVM file with a real libNVVM, ptxas and nvJitLink run.
  • Widened the targets to the full Blackwell family list; the f-only list refused --arch sm_120a and sent an arch-less run on a 5090 to sm_100f.
  • Added the rule that PTX 9.1 needs LLVM 23: llc-22 silently emitted .version 8.8 PTX that failed at load. The examples CI lane moves to LLVM 23 for the same reason.
  • Merged main, regenerated, added the negative admission tests and the four forms to generated_intrinsics_blackwell.

At 89e24e263, the catalog gates, the crate tests, the guards, the Blackwell example on sm_120a with memcheck clean, and all CI checks passed.

@nihalpasham
nihalpasham merged commit e5082e6 into NVIDIA:main Oct 19, 2026
64 checks passed
@YeonwooSung
YeonwooSung deleted the feat/bf16x2-to-fp8x2-pack branch October 20, 2026 03:45
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

build-related Build scripts, toolchain discovery, or compilation setup cuda-feature A CUDA hardware/toolkit feature (TMA, tcgen05, WGMMA) cuda-oxide SIMT programming model: rustc backend, cargo oxide, cuda-device, cuda-host, book device-apis User-facing device-side and kernel-authoring APIs (cuda-device, cuda-macros) enhancement New feature or request intrinsics Device intrinsics and libdevice math mappings

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants