You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
{{ message }}
Repository navigation
feat(rocke): add logical TF32 and native gfx942 atoms - #12612
TF32 needs a logical identity independent of its storage carrier so catalog queries can select native XF32 instructions without changing ordinary FP32 selection. This adds tf32 (xf32 alias), represented as I32 words, and the gfx942 16x16x8 and 32x32x4 native atoms in both engines. Lowering bitcasts the carrier to the intrinsic's float-vector ABI.
Raw FP32 payload wrapping is separate from explicit round-to-nearest-even preparation. The reusable numerical example demonstrates that raw inputs are accepted on gfx942 and match truncation for the tested corpus; RNE-prepared inputs differ at rounding boundaries. For example, multiplying input bits 3f801001 by one produces 3f800000 through raw XF32 and 3f802000 after RNE preparation. I32 transport preserves the raw result.
This reuses the fragment packing/loading helpers now present in the develop base. It includes dtype/serialization/transport support, operand and target guards, a mirrored probe, documentation, and ten new LLVM golden records. Generic scalar and vector TF32 arithmetic is rejected; payload transport, bitcasts, and selection remain supported. Direct global vector loads accept 2, 3, 4, or 8 TF32 values; stores accept 1, 2, 4, or 8, using I32 alignment rules. Payload width and address alignment are independent. Both LLVM store lowerers honor the explicit alignment, and HIP uses an exact-byte copy for underaligned stores. Ordinary MMA atoms also reject TF32 result types. Existing golden records are unchanged. Production GEMM dispatch and cross-target emulation are outside this change.
Architecture and input preparation:
flowchart TD
F["FP32 input payloads"] --> L["Existing fragment packing and loading helpers"]
I["Prepared I32 input words"] --> L
L --> W["Bitcast: preserve raw payload"]
L -->|FP32 values| R["Optional FP32 to TF32 RNE recipe"]
W --> T["Logical A/B: two TF32 values per lane"]
R --> T
Q["Catalog query: tf32, tf32, fp32"] --> A["gfx942 XF32 atom: 16x16x8 or 32x32x4"]
T --> IR["Typed tile.mma IR, serialization and validation"]
A --> IR
C["FP32 accumulator C"] --> IR
IR --> PY["Python lowering"]
IR --> CPP["C++ lowering"]
PY -->|A and B| S["I32 carrier: LLVM 2 x i32 / HIP i32x2"]
CPP -->|A and B| S
PY -->|C| ACC["FP32 accumulator vector"]
CPP -->|C| ACC
S --> ABI["A/B bitcast: LLVM 2 x float / HIP f32x2"]
ABI --> M["Native gfx942 XF32 MFMA"]
ACC --> M
M --> D["FP32 result D"]
Loading
Logical TF32 drives catalog selection and operand validation. The I32 carrier preserves A/B payload bits; the float-vector ABI cast does not round them. C and D remain FP32. RNE preparation expands to existing integer IR operations, with no new serialized opcode. The two LLVM emitters are checked for byte identity.
Slurm job 68011208 completed with exit status 0:0 on a real gfx942. The Python source and freshly built C++ bindings came from commit 4f18490d8ce6314ef06014a5f951a38312caa411; the embedded source archive was verified before extraction (SHA-256 625c9e203ff7a35e6ac9eb5907345272a26434db05b7a29f5826f3bd5b9bd2a1).
Both engines × both native shapes (16x16x8, 32x32x4) × six variants passed with normal Python, then passed again with python -O: 24 + 24 successful variant runs, each with 2,156 batched atoms. Variants were raw FP32 payloads, I32 transport, GPU RNE preparation, host RNE preparation, host masking, and ordinary FP32 MFMA.
The unmodified numerical harness checked actual prepared A/B words, exact single-product boundaries (including subnormals and special values), dense inputs, cancellation, and nonzero accumulators against independent NumPy references. Existing tolerances were unchanged. Python/C++ builder and LLVM-lowering parity, cross-variant equality, and cross-engine output equality passed. Input bits 3f801001 still produced 3f800000 through raw/I32 transport and 3f802000 after RNE preparation.
Environment: ubuntu-24 + rocm/7.14.0, HIP runtime 7.14.60850, LLVM23, Python 3.12.3, NumPy 2.4.6, and freshly built native bindings using GCC 13.2.0. HIP, ROCr, and COMGR were loaded from the same ROCm module. Kernels were compiled from LLVM IR through COMGR and launched through HIP; PyTorch was not used.
All 48 generated code objects contained the expected MFMA instructions, with XF32 present for TF32 variants and absent for the ordinary FP32 control. Corresponding LLVM IR and code objects were byte-identical between the normal and optimized Python runs. Public TF32 export identity also passed on the compute node.
Two earlier attempts (68011174, 68011200) on another gfx942 node aborted with SIGBUS at hipStreamSynchronize before numerical comparison. Their first compiled object is byte-identical to the successful run's object. These attempts provide no passing numerical result, and the cause of that node's failure remains undiagnosed.
Public Python API validation (4f18490d8ce6): from rocke import TF32 and from rocke.core import TF32 both expose the canonical IR type and include it in __all__. All 152 focused TF32 tests and full PR-range pre-commit hooks passed. This export-only follow-up does not change generated kernels; current-head GPU validation is recorded above.
Numerical harness tightening (1c75ffdf3c01):
All eight parity/numerical checks now use explicit conditional exceptions, so validation remains active with python -O and PYTHONOPTIMIZE=1.
178 focused host tests passed, including 27 subprocess cases covering success and all eight failure paths under normal Python, -O, and PYTHONOPTIMIZE=1. Injected failures verify nonzero exit status, results.json failure status, and diagnostic context. Compiler/GPU calls are mocked in these new tests; they establish validation-control behavior, not GPU numerical correctness.
Ruff and full PR-range pre-commit hooks passed. Kernel builders, emitters, goldens, and numerical tolerances are unchanged. The full all-family byte-identity gate was not rerun for this host-only change; current-head GPU validation is recorded above.
Validation after merging develop and tightening the API (a09dd3378f2b):
151 focused host tests passed with freshly built native bindings. Coverage includes serialized TF32 vector arithmetic rejection, scalar/vector TF32 results on ordinary MMA atoms, positive payload transport, direct vector loads, and store alignment for TF32/I32/F32. Default, weaker, stronger, and absent legacy alignment attributes are covered; zero, negative, and non-power-of-two serialized store alignments are rejected consistently by LLVM and HIP, with Python/C++ LLVM parity for LLVM20/22/23.
All 14 native CTests passed, including native HIP rejection and aligned/underaligned store lowering. At the preceding tightening revision (75508e81e332), eight direct TF32 store kernels assembled as LLVM IR and compiled as HIP device code for gfx942. The four-element, four-byte-aligned HIP case still emits global_store_dwordx4; the alignment fix preserves the wide transfer in that tested case. The prior eight direct vector-load compilation checks also passed.
All 71 families / 816 sampled configurations were exercised in each LLVM20/22/23 byte-identity run using local system Python. TF32 passed in every flavor. The remaining gfx942_attention_tiled_3d config 3 mismatch is inherited from develop: it was reproduced on the exact base with identical per-engine output hashes (Python includes _fnuz in the kernel name; C++ does not).
Combined TF32, dtype, storage, serialization, and MMA host coverage at 75508e81e332: 358 passed, 39 skipped, and four failures. The four storage failures are inherited from develop and concern the removed ScalePacking.llvm_type property.
Final local architecture/API review closed the vector arithmetic, vector-load admission, store alignment and its validation, and ordinary-MMA result-type gaps using existing mirrored helpers. The canonical dtype reference distinguishes load/store widths and address alignment. Repository pre-commit hooks passed over the full PR range.
Earlier implementation validation (before the develop merge and this tightening):
Slurm gfx942, ROCm 7.14: both engines × both shapes × six variants passed, with 2,156 batched atoms per variant. Checked prepared input words, exact single-product boundaries (including subnormals and special values), dense inputs, cancellation, and nonzero accumulators against independent references. The 24 programs at that revision matched the hardware-tested artifacts and contained the expected instructions.
At that revision, all ten probe modes compiled through COMGR and HIP on LLVM20/22/23, and Python/C++ HIP emission matched byte-for-byte.
Windows validation was not available; LDS/loop transport has host lowering/parity coverage rather than a separate GPU LDS round-trip test.
Reproduce the numerical example with a matching ROCm compiler/runtime and built native engine on PYTHONPATH:
python -m rocke.examples.gfx942.tf32_numerics --backend both --shape all --output-dir resultspython -O -m rocke.examples.gfx942.tf32_numerics --backend both --shape all --output-dir results-optimized
Keep empty descriptors valid while rejecting fragment loads before IR emission in both engines. Cover empty shapes and nonempty controls in Python and native tests.
Validate complete static fragment displacements before emitting i32 offsets in both engines. Resolve CTest readiness from registered commands and cover multi-config execution and failure propagation.
Size temporary load-pointer storage by addressable units so padded slots
cannot overrun the array. Preserve IR order and add mirrored transport
coverage for a layout requiring more loads than logical elements.
Simplify scale-width queries while retaining the shared packing adapter,
and clarify the contiguous-K matrix layout contract.
The new logical type is not re-exported from either rocke.core or the top-level rocke package, unlike the existing public scalar types such as F32. As a result, the normal authoring import from rocke import TF32 (and from rocke.core import TF32) fails even though TF32 is now part of the public IR API. Please add TF32 to both packages' imports and __all__ lists.
Raw FP32→TF32 truncates (HW drops low 13 mantissa bits); RNE is opt-in, so future GEMM dispatch wrapping raw FP32 is silently less accurate than framework "TF32"
2
🟡
PR validation
GPU pass came from one gfx942 node; two others aborted undiagnosed SIGBUS at hipStreamSynchronize on a byte-identical object — note before broad reliance
V1
🟢
ir_arith.cpp cvt_f32_to_tf32
RNE recipe verified correct: bias 0x0FFF+odd, mask 0xFFFFE000, Inf passthrough, sNaN→qNaN with no NaN→Inf collapse, ±0/subnormals preserved
V2
🟢
tf32.py/tf32_internal.h
TF32 guards now check results on XF32 and non-XF32 paths; generic arith on i32 carrier blocked via allowlist — symmetric both engines
3
🟢
regression status
"Red" is only inherited failures (gfx942_attention_tiled_3d cfg3 + removed ScalePacking.llvm_type), not caused here
@cgmillette
RNE vs truncation is a design decision. It needs a bit of further study what is the common framework convention and we may flip the default if needed. Both code paths are enabled so it looks like a matter of framework integration rather than compiler. I think it shouldn't be blocking
SIGBUS diagnostics need stress testing. I don't think this is blocking either
…port
The installed Linux test artifact cannot run two TF32 tests introduced
in #12612. Recipe replay attempts to configure a native source build
from the installed bin directory, which has no CMake source tree. The
gfx942 numerical test requests both backends and imports rocke_engine,
which is absent from the installed artifact.
Add direct pytest skip decorators to the ten recipe-replay cases and
the one gfx942 numerical case pending fixes to their artifact support.
Keep the remaining TF32 tests enabled and include the originating PR
in each skip reason. These temporary skips also apply to local runs.
Validation: targeted host tests passed (34 passed, 127 skipped), and
repository pre-commit checks passed. The gfx942 Enroot/Slurm run honors
all eleven skips and passes both SDPA suites; the broader platform
suite still has six separate FP4 alias test failures.
Resolves conflicts with #12612 (logical TF32 + native gfx942 atoms), which
taught the plain HIP global_store_vN to honour `align` (memcpy fallback when
under-aligned) in both engines. The store now mirrors the load exactly:
under-aligned or non-power-of-two payloads take the memcpy path, and a
STREAMING hint there is refused (NotImplementedError / ROCKE_ERR_NOTIMPL),
since the HIP backend does not yet lower nontemporal through memcpy.
ir.py: global_store_vN docstring keeps develop's alignment text plus the
temporal_hint paragraph. Tests updated to the memcpy-path wording.
The nontemporal_emit parity family (4 copy-kernel configs: streaming /
default loads and stores on gfx942 / gfx950) had cross-engine parity and
substring checks but no golden entry, so an identical regression in both
lowerers would go unnoticed. Record its Python reference shas under both
golden keys (`ll` and `ir_canonical`), as #12612 did for tf32_mma.
Recorded with the family GREEN in both modes (C++ == Python). Only the
`nontemporal` entries are added (12 lines, 0 changed): `--record-golden`
rewrites the whole key from the current run, and most other families do
not match the committed golden on an llvm20 host even on develop, so
re-recording everything would rewrite unrelated hashes. tf32_mma, the
latest blessed family, matches on the same host.
The installed rocKE suite failed eleven TF32 cases introduced by #12612:
ten recipe cases attempted a native source build from the installed
`bin` directory, and the gfx942 numerical test requested `rocke_engine`,
which the artifact does not contain. #12718 temporarily skipped these
tests in every environment.
Restore the tests by installing the existing native recipe replay CLI
under `tests/portable_ir/`, which TheRock already captures in the test
artifact. Compare native CBOR replay with Python LLVM lowering for both
atom shapes, all preparation modes, and LLVM 20/22/23. Installed CTest
supplies a relocatable executable path through the existing
`ROCKE_REPLAY_CLI` setting for all native CLI replay tests. TF32 keeps
its module-local fixture and uses the same setting. This also enables
the four existing portable-IR CLI tests in installed artifacts. An unset
or missing replay executable skips the native lane, matching the
existing portable-IR tests; execution errors and IR mismatches remain
failures.
Run all 12 Python numerical variants against the independent NumPy
references when the native binding is absent. When the binding is
available, require all 24 variants, native builder/lowerer byte
identity, and backend numerical parity; a broken binding remains a
failure. The standalone numerical CLI retains its strict both-backend
default. Document general native replay rules in `TESTING.md`: use
prebuilt artifacts, skip absent CLIs, avoid implicit source builds,
isolate fixtures, and validate relocated installs.
Fixes#12718
ISSUE ID : AICK-2514
## Test plan
- [x] Reproduce all ten original recipe source-build failures in an
isolated installed layout, then pass all ten cases with matching newly
built artifacts.
- [x] Fresh native build and focused TF32 core/wrapper suite before
removing three wrapper tests: **155 passed**. Existing numerical
fault-injection coverage retained.
- [x] Native TF32, DOM decoder, and recipe VM CTests: **3 passed**.
- [x] TF32 core and numerical-validation source suite without binding:
**63 passed, 116 capability skips**.
- [x] Full relocated installed `rocke_pytest` with the shared replay
setting: **1,727 passed, 321 skipped, 1,632 subtests passed**.
- [x] Matching platform/library wheels installed with `--no-deps
--reinstall`; CTest discovery confirmed the selected entry and installed
import root.
- [x] Single-setting replay check: all 14 TF32/portable-IR CLI cases
pass in source, fresh installed, and relocated layouts; unset or missing
CLI gives 14 skips and a broken executable gives 14 failures. Provider
and standalone CTest paths verified.
- [x] Repository pre-commit hooks over the PR range and whitespace
checks.
- [ ] Real gfx942 numerical execution and native Windows validation:
deferred to CI; no new GPU result is claimed.
The local installed replay used the platform's provider-style
installation and generated provider CTest entries, rather than the
complete provider superbuild. Tools were Python 3.12.12, NumPy 2.5.3,
pytest 9.1.1, pytest-timeout 2.4.0, ml-dtypes 0.6.0, and ROCm 7.1 Clang
20; these differ from the original CI runtime. Kernel emitters and
numerical tolerances are unchanged.
An initial full run used `/usr/bin/llc` and failed four unrelated
debug-info object tests because that compiler does not recognize gfx950.
The reported passing run selects ROCm LLVM via `ROCKE_TEST_LLVM_BIN` and
`PATH`.
…#13588)
# feat(rocke): emulate logical TF32 MMA on gfx950 with three BF16 MFMAs
Branch: `users/congma13/xf32-bf16x3-emulation-v2`
ISSUE ID : AICK-2208
---
# Summary — For human reviewers
## The short version
gfx942 has a hardware matrix instruction that multiplies fp32 values at
reduced
precision (a 19-bit format; the instruction names call it XF32, and it
is commonly
called TF32). gfx950 does not. #12612 added a logical `tf32` type and
the two native gfx942
atoms, so a kernel can ask the arch catalog for "a TF32 matrix op" and
get an
answer on gfx942 — and nothing at all on gfx950.
This PR adds two gfx950 catalog atoms, `mfma_f32_16x16x32_xf32` and
`mfma_f32_32x32x16_xf32`, that answer the same query. Each one lowers to
three
BF16 matrix instructions over a split of the fp32 operands, in both the
Python
and the C++ engine, with byte-identical output. They take the **same
operand
type** as the native atoms, so one kernel body can reach either.
This is the whole gfx950 capability. Calling it from production
dispatch, and an
accuracy-driven option on gfx942, are separate changes (see "Reviewing
this").
## Why the operand type has to match the native atoms
The first version of this branch, and the obvious design, took ordinary
fp32
vectors. It would have worked, and nothing would have objected: the TF32
verifier from #12612 only rejects TF32-typed operands on atoms it does
not know,
so an fp32 vector slips straight through.
The cost only shows up when someone writes a kernel. Asked for
`(tf32, tf32, fp32)`, the catalog would hand back an atom that wants
`vec<tf32x2>` on gfx942 and one that wants `vec<f32x8>` on gfx950. The
one thing
emulation is for — write the kernel once, let the catalog pick native or
emulated — would not exist. Handing the gfx950 atom the TF32 type the
query
promises would have been rejected by that same verifier.
```mermaid
flowchart LR
Q["kernel asks for<br/>tf32, tf32 → fp32"] -->|gfx942| N["native XF32 atom"]
Q -->|gfx950| E["emulated atom"]
N -->|operands| T["vec of tf32<br/>(I32 carrier)"]
E -->|operands| T
E -->|lowering| L["bitcast, split,<br/>3 BF16 MFMAs"]
```
Both atoms take the same operand type; only the lengths differ (two
values per
lane natively, eight on gfx950). The accumulator is an fp32 vector on
both.
### Emulating TF32 does not mean rounding to TF32
The second surprise. A native XF32 instruction ignores the low 13 bits
of each
operand. The emulation could have copied that, and would have matched
the hardware
more closely. It does not: it reinterprets the whole 32-bit payload as
fp32 and splits that, so every bit the caller supplied reaches the
result. The
base's own boundary example is an operand `0x3f801001` multiplied by
1.0:
| path | result |
|---|---|
| native atom, raw payload | `0x3f800000` — low bits dropped |
| native atom, after explicit round-to-nearest-even preparation |
`0x3f802000` — rounded to 10 mantissa bits |
| BF16x3 emulation | `0x3f801000` — the payload, to within one fp32 ulp
|
All three are right for what each is. The emulation's contract is
therefore:
*track the exact product of the payloads it is given, to within the
split's own
error; not bit-reproducible against the native atom across targets.* On
raw fp32
payloads that is more accurate than the native atom. On payloads already
rounded
to TF32 precision the native atom is exact, and the emulation stays
within `2^-18`
(relative, per product) of it — the bound on the one term it omits.
## How it works
Each emulated `tile.mma` is expanded at lowering, before the
instruction-set
dispatch:
1. **View the payload as fp32.** A `bitcast` of the I32 carrier — never
a
rounding.
2. **Split each operand** into `big = bf16(x)` and `small = bf16(x -
fp32(big))`.
Both narrowings are `fptrunc`, which LLVM defines as
round-to-nearest-even and
which the AMDGPU backend selects as the packed convert on gfx950.
3. **Issue three BF16 MFMAs** threading one fp32 accumulator, in the
order
`small*big`, `big*small`, `big*big`. The order is part of the answer:
fp32
accumulation is not associative, so any other order gives different
bits. The
fourth product, `small*small`, is omitted; it is the same size as the
split's own
representation error.
The decisions that are not optional:
- **The verifier reads the catalog.** #12612's TF32 contract is a
two-entry table
keyed by op_id (Python) and a pair of `strcmp`s (C++), plus five
hand-written
`gfx942` guards. Adding the emulated ids to it would have merged two
properties it
holds in one place: *which atoms take TF32 operands* and *which atoms
are native
gfx942 instructions*. Only native atoms existed, so they coincided. They
no longer
do, and the consequences are concrete: the HIP emitter builds
`__builtin_amdgcn_<op_id>` from that table, so an emulated id there
would emit a
builtin that does not exist, and the `gfx942` guard would reject the
emulated atom
on gfx950. So the emulated atoms get their own catalog-driven branch —
operand and
accumulator lengths come from the row — and the native table and guards
are left
exactly as they were.
- **The catalog cannot over-claim.** An emulated row's physical atom
must be a native
row of the *same* arch, checked when the catalog loads. The first
version on this
branch asked only whether the atom had fragment metadata anywhere, which
would have
let a row claim support the compiler then rejects at codegen ("Cannot
select"). The
recipe also fixes the operand dtype. The C catalog is static data, so
its twin is a
test that walks each emulated row's own arch.
- **The native-preferred rule is pinned, not changed.** Selection drops
an emulated
atom only when a native atom of the *same* `(m, n, k)` is listed. A
native atom with
a different K does not shadow it, and `select_largest_k` would then
return the
emulated atom — three MFMAs instead of one. This was reproduced by
running the real
catalog code on a synthetic native-K=8 plus emulated-K=16 or K=32
catalog. Nothing
triggers it, because no arch lists both; a test pins that, and the
rule's docstring
says where to re-key it.
- **HIP rejects an emulated atom by name.** The expansion exists only in
the LLVM
lowerer. Before, an emulated op reached the HIP lowerer's generic
dispatcher and
failed with "no HIP lowering for op `tile.<op_id>`", which, with a
hardware-shaped
name, reads as a missing hardware atom.
- **gfx942 is deliberately untouched.** The two atoms here expand into
the K-packed
bf16 MFMAs, which exist only on gfx950; the compiler-probed intrinsic
table records
them as unavailable on gfx942. A gfx942 emulation would be two different
atoms
(`16x16x16`, `32x32x8`), would hit the selection rule above, and would
put two
numerics behind one dtype on one arch.
- **Two earlier designs are removed.** The branch first modelled the
dtype on CK as an
`xf32` selection tag over fp32 operands, and added an optional MMA
result-dtype field
for it. #12612's `tf32` supersedes the first; the second never had a
user. Neither is
in the net diff; the history still shows them.
## What this found
- This branch's first-version catalog check had the hole described
above; it is closed
here.
- Before the parity family below, **no kernel family in the
byte-identity gate selected
an emulated atom**, so the gate could not have compared the two engines
on the
expansion at all.
- On the merged base the byte-identity gate is green outright. The
single divergence
earlier versions of this work carried (`gfx942_attention_tiled_3d`, one
config)
no longer reproduces; nothing here touches it.
- The L5 golden anchor (`llvm_gfx_all.json`) is mostly stale on
`develop`: at the
flavor of record, 11 of its 73 `ll` families match the tree entirely.
That is not
this PR's doing and it is left alone; see the AI-reviewer notes.
## How this was tested
The numerics are pinned three ways: a host model of the recipe (no GPU),
an
on-device harness that compares the emulated `tile.mma` bit for bit
against a
hand-written kernel issuing the same three BF16 MFMAs, and an absolute
check
against `K*alpha*beta` with a single-term control that must miss. The
two engines
are compared by a new parity family in the byte-identity gate, which was
itself
shown not to be vacuous: swapping two terms in the C++ expansion alone,
in a scratch
copy, turns both of its configs into DRIFT.
### Test plan
- [x] `tests/core` with the `rocke_engine` module built, plus
`tests/test_rocke_ci_static.py` and `tests/test_rocke.py`: 1635 passed,
79
skipped, 2 failed. The two failures also fail on unmodified `develop` in
the same
environment: `test_arch_domain_artifact.py` — the `llvm22`
intrinsic-domain
artifact lacks 26 recently added cluster and prefetch keys, and its
"regenerating this host's column is a no-op" check follows from that.
- [x] C++ tests under `ctest`: `rocke_tf32`, `rocke_dtypes`,
`rocke_mma_catalog_lookup`, `rocke_xf32_bf16x3_lowering`,
`rocke_mma_frag_ssot`.
- [x] COMGR compile gate: both emulated atoms compile for gfx950 to a
code object with
exactly three `v_mfma_*_bf16` and the packed convert.
- [x] `tools/check_byte_identity.py` at all three flavors, run
sequentially: every
family GREEN (77 per flavor), no DRIFT, including the base's native TF32
family
and the new parity family.
- [x] The `xf32_bf16x3_mma` parity family re-diffed Python-vs-C++ on the
final archive
at all three flavors, in both `ll` and canonical `ir` mode.
- [x] Negative control for that family (scratch copy, C++ expansion
only): DRIFT on
both configs at every flavor in `ll`; GREEN on the real archive.
- [x] On-device harness on gfx950 — Python lowerer, C++ engine with the
Python
fallback disabled, and both engines cross-checked: bitwise equal to the
hand-written BF16x3 reference at both atoms; accuracy inside tolerance;
the
single-term control outside its floor.
- [x] On-device fault injection (term order rotated; residual terms
dropped): both
rejected with exit status 2, unmodified control exits 0.
- [x] `tests/core/test_xf32_bf16x3_device.py` under
`ROCKE_REQUIRE_GFX950=1` on the same
machine, including with `HIP_VISIBLE_DEVICES` set.
- [x] Representative-IR golden
(`tests/golden/rocke_representative_ir_sha256.json`)
checks clean at all three flavors, unchanged.
- [x] Full `--check-golden` of the L5 anchor at the flavor of record,
both modes: the new
family and the base's native TF32 family match; older entries are stale
on the
base and untouched.
- [x] `pre-commit run --from-ref origin/develop --to-ref HEAD` passes
(black,
clang-format, bandit, cmake-lint, whitespace hooks).
- [ ] **Deferred:** the full three-flavor byte-identity gate was run
before the last
five commits. Those touch no emission path — a verifier-lookup guard
that raises
on malformed catalog rows, a HIP-lowerer message, tests, a golden JSON,
docs, and
one comment — and the parity family, the only family whose emission this
PR
changes, was re-diffed on the final archive. A re-run on the final head
is the
remaining step.
- [ ] **Not run:** the base's native TF32 numerics tests (need gfx942
hardware; that
code path is untouched).
- [ ] **Not run:** a node where the gfx950 is not HIP device 0. The
probe's index
selection was verified with faked probe output, not on silicon.
- [ ] **Not run:** CI GPU lanes and other host platforms.
---------
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
github-actions
Bot
added
documentation
project: hip-kernel-provider
rocke: platform
ISSUE ID : AICK-1714
TF32 needs a logical identity independent of its storage carrier so catalog queries can select native XF32 instructions without changing ordinary FP32 selection. This adds
tf32(xf32alias), represented as I32 words, and the gfx94216x16x8and32x32x4native atoms in both engines. Lowering bitcasts the carrier to the intrinsic's float-vector ABI.Raw FP32 payload wrapping is separate from explicit round-to-nearest-even preparation. The reusable numerical example demonstrates that raw inputs are accepted on gfx942 and match truncation for the tested corpus; RNE-prepared inputs differ at rounding boundaries. For example, multiplying input bits
3f801001by one produces3f800000through raw XF32 and3f802000after RNE preparation. I32 transport preserves the raw result.This reuses the fragment packing/loading helpers now present in the develop base. It includes dtype/serialization/transport support, operand and target guards, a mirrored probe, documentation, and ten new LLVM golden records. Generic scalar and vector TF32 arithmetic is rejected; payload transport, bitcasts, and selection remain supported. Direct global vector loads accept 2, 3, 4, or 8 TF32 values; stores accept 1, 2, 4, or 8, using I32 alignment rules. Payload width and address alignment are independent. Both LLVM store lowerers honor the explicit alignment, and HIP uses an exact-byte copy for underaligned stores. Ordinary MMA atoms also reject TF32 result types. Existing golden records are unchanged. Production GEMM dispatch and cross-target emulation are outside this change.
Architecture and input preparation:
flowchart TD F["FP32 input payloads"] --> L["Existing fragment packing and loading helpers"] I["Prepared I32 input words"] --> L L --> W["Bitcast: preserve raw payload"] L -->|FP32 values| R["Optional FP32 to TF32 RNE recipe"] W --> T["Logical A/B: two TF32 values per lane"] R --> T Q["Catalog query: tf32, tf32, fp32"] --> A["gfx942 XF32 atom: 16x16x8 or 32x32x4"] T --> IR["Typed tile.mma IR, serialization and validation"] A --> IR C["FP32 accumulator C"] --> IR IR --> PY["Python lowering"] IR --> CPP["C++ lowering"] PY -->|A and B| S["I32 carrier: LLVM 2 x i32 / HIP i32x2"] CPP -->|A and B| S PY -->|C| ACC["FP32 accumulator vector"] CPP -->|C| ACC S --> ABI["A/B bitcast: LLVM 2 x float / HIP f32x2"] ABI --> M["Native gfx942 XF32 MFMA"] ACC --> M M --> D["FP32 result D"]Logical TF32 drives catalog selection and operand validation. The I32 carrier preserves A/B payload bits; the float-vector ABI cast does not round them. C and D remain FP32. RNE preparation expands to existing integer IR operations, with no new serialized opcode. The two LLVM emitters are checked for byte identity.
Current-head gfx942 hardware validation (
4f18490d8ce6, 2026-09-30):68011208completed with exit status0:0on a real gfx942. The Python source and freshly built C++ bindings came from commit4f18490d8ce6314ef06014a5f951a38312caa411; the embedded source archive was verified before extraction (SHA-256625c9e203ff7a35e6ac9eb5907345272a26434db05b7a29f5826f3bd5b9bd2a1).16x16x8,32x32x4) × six variants passed with normal Python, then passed again withpython -O: 24 + 24 successful variant runs, each with 2,156 batched atoms. Variants were raw FP32 payloads, I32 transport, GPU RNE preparation, host RNE preparation, host masking, and ordinary FP32 MFMA.3f801001still produced3f800000through raw/I32 transport and3f802000after RNE preparation.ubuntu-24+rocm/7.14.0, HIP runtime7.14.60850, LLVM23, Python3.12.3, NumPy2.4.6, and freshly built native bindings using GCC13.2.0. HIP, ROCr, and COMGR were loaded from the same ROCm module. Kernels were compiled from LLVM IR through COMGR and launched through HIP; PyTorch was not used.68011174,68011200) on another gfx942 node aborted with SIGBUS athipStreamSynchronizebefore numerical comparison. Their first compiled object is byte-identical to the successful run's object. These attempts provide no passing numerical result, and the cause of that node's failure remains undiagnosed.Public Python API validation (
4f18490d8ce6):from rocke import TF32andfrom rocke.core import TF32both expose the canonical IR type and include it in__all__. All 152 focused TF32 tests and full PR-range pre-commit hooks passed. This export-only follow-up does not change generated kernels; current-head GPU validation is recorded above.Numerical harness tightening (
1c75ffdf3c01):python -OandPYTHONOPTIMIZE=1.-O, andPYTHONOPTIMIZE=1. Injected failures verify nonzero exit status,results.jsonfailure status, and diagnostic context. Compiler/GPU calls are mocked in these new tests; they establish validation-control behavior, not GPU numerical correctness.Validation after merging develop and tightening the API (
a09dd3378f2b):75508e81e332), eight direct TF32 store kernels assembled as LLVM IR and compiled as HIP device code for gfx942. The four-element, four-byte-aligned HIP case still emitsglobal_store_dwordx4; the alignment fix preserves the wide transfer in that tested case. The prior eight direct vector-load compilation checks also passed.gfx942_attention_tiled_3dconfig 3 mismatch is inherited from develop: it was reproduced on the exact base with identical per-engine output hashes (Python includes_fnuzin the kernel name; C++ does not).75508e81e332: 358 passed, 39 skipped, and four failures. The four storage failures are inherited from develop and concern the removedScalePacking.llvm_typeproperty.Earlier implementation validation (before the develop merge and this tightening):
Reproduce the numerical example with a matching ROCm compiler/runtime and built native engine on
PYTHONPATH: