feat(rocke): add i4 and unsigned types, explicit f64 conversions, bit-width and packing contracts, and MMA operand typing by cgmillette · Pull Request #13434 · ROCm/rocm-libraries

cgmillette · 2026-10-09T23:10:18Z

Motivation

rocKE needs a few more data types and a few firmer core contracts before the tiling layer can work with mixed and low-precision MMA instructions. Today:

there is no logical i4 and no unsigned 8- or 4-bit type;
iu8 / iu4 are their own dtypes, although signedness is really an instruction operand;
converting between f32 and f64 has no explicit op;
a width-changing bitcast is only caught by clang;
nothing in core says how sub-byte values pack into 32-bit registers;
some MMA lowerings emit IR that LLVM rejects when an operand arrives in another shape, for example a <1 x float> fragment for the f32 MFMA, which the tiling layer builds;
MMA results are typed "i32 else f32";
TF32 MMA availability is a hardcoded gfx942 check.

This PR adds the types and fixes those contracts in both engines. Existing kernels are not affected.

f64 itself comes from #13141, which adds it to both engines. This PR builds on #13141 and does not add its own f64 support.

ISSUE ID : AICK-2568

Also completes AICK-1713 together with #13141: #13141 delivered the f64 type, memory and arithmetic, and this PR adds the f32 <-> f64 conversions that ticket also asks for.

Technical Details

Types and conversions

I4 and U4 are packed-only logical types: their values live packed in 32-bit words, and the IR builder refuses to create an i4 / u4 scalar or vector value, with the same message in both engines. The 8- and 4-bit integer dtypes are i8 / i4; iu8 / iu4 are no longer dtypes (they stay only in the RDNA integer WMMA instruction names).
U8 is an unsigned 8-bit arithmetic type: division, remainder, ordered compares, int <-> float conversion, widening, min / max and right shift use the unsigned forms in both engines.
cvt_f32_to_f64 / cvt_f64_to_f32 are explicit conversion ops, each one instruction, with their own opcodes in both engines.
LDS loads and stores accept i16.

Core contracts

type_bits(t) is the one place that answers how many bits a value of a type occupies; pointers are sized by address space. bitcast / vec_bitcast refuse a width change and refuse pointers.
register_packing / packed_word_range in core/storage.py describe how sub-byte values pack densely into 32-bit registers. They build on the existing BitPacking / FragmentPacking.
test_dtypes.py has a hand-written width table covering every IR scalar type. dsl_docs/development/extending.md section 5.4 is the single recipe for adding a dtype.

MMA lowering

Operands are passed as the instruction's register type (_operand_as / _mma_operand).
- A same-width operand of the instruction's own format in another shape is reinterpreted with one bitcast, for example <1 x float>, or packed <16 x i8> / <32 x i8>.
- Any other number format raises instead of being silently reinterpreted.
Results are typed by the catalog's accumulator dtype.
TF32 MMA availability is read from the target's MMA catalog.

Both engines

Every change above is made in the Python engine and the C++ engine, and both emit the same LLVM IR.

Tooling

The IR-validity checker looks in the installed ROCm before PATH.
rocke.tests.ir_validity gives tests a compile-and-link check.
rocke.tests.gpu runs numeric tests only on the arches they name, and any failure there is loud.

Status

Built on feat(rocke): add FP64 support for elementwise and reduce #13141, which provides f64 in both engines; this PR adds no f64 support of its own.
Every change is mirrored in the C++ engine; both engines emit identical LLVM IR.
New opcodes are appended to the public opcode enum, so no existing opcode value changes; a test pins every pre-existing value.

Test Plan

New unit tests for the new types, conversions, bitcast width, register packing and MMA operand typing, with additions to the existing dtype, storage, MMA table, TF32, store_fragment, IR-validity and arch-catalog tests.
Tests that build each change with both engines and compare the IR.
Lowered IR in the new tests is compiled and linked for its target: gfx90a, gfx942, gfx950 and gfx1151 with the host ROCm clang, and gfx1250 with ROCm 10's clang.
Numeric GPU tests for the conversions and LDS round trips run on gfx90a.
Byte identity between the C++ and Python engines is checked for every parity kernel family at every LLVM flavor.
The full platform, tiling and library test suites are run on gfx90a.
Blast-radius census: every MMA call across all parity kernel families and the full suite is recorded, comparing operand and result types before and after.
black, ruff and pre-commit.

Test Result

Byte identity: green at llvm20, llvm22 and llvm23 for every kernel family. Existing kernels emit identical IR.
Each of the 11 commits builds and passes its own tests on its own.
Strict both-engines mode (ROCKE_BACKEND=both ROCKE_CPP_STRICT=1) over the PR's test files: all pass.
GPU tests on gfx90a (f64 edge values and conversions, u8 arithmetic against numpy uint8): bit-exact.
Full suites on gfx90a: no new failures. The remaining failures also fail on the base commit and are unrelated (environment-dependent tests).
C++ tests (ctest): all pass.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

🤖 Generated with Claude Code

therock-pr-bot · 2026-10-09T23:15:04Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

therock-pr-bot · 2026-10-09T23:15:09Z

🎉 All checks passed! This PR is ready for review.

Copilot

🟡 Changes recommended

The default C++ engine remains inconsistent with new Python emission paths, and F64 LDS allocations can be under-aligned.

5 open findings

Mirror canonical operand rows in C++ tables · New
Add new public scalar types to the C++ registry · New
Mirror operand bitcast handling in the C++ lowerer · New
Align f64 allocations to 8 bytes · New
Preserve conversion edge cases in GPU assertions · New

What changed in this PR

Adds Python-side dtype, packing, fragment-safety, and MMA typing contracts for low-precision and mixed-type kernels.

Changes:

Adds F64, I4, U4, U8, type-width, bitcast, and register-packing support.
Strengthens fragment/store contracts and MMA operand/result typing.
Adds IR-validity and GPU test infrastructure with supporting documentation.

File	Description
`platform/tools/check_ir_validity.py`	Searches installed ROCm toolchains before PATH.
`platform/tests/core/test_tf32.py`	Tests catalog-driven TF32 rejection.
`platform/tests/core/test_storage.py`	Tests register packing contracts.
`platform/tests/core/test_mma_operand_register_type.py`	Tests MMA operand reinterpretation and validation.
`platform/tests/core/test_mma_frag_tables.py`	Tests catalog-derived accumulator types.
`platform/tests/core/test_f64.py`	Tests F64 lowering, memory, and execution.
`platform/tests/core/test_dtypes.py`	Tests new types and scalar widths.
`platform/tests/core/test_bitcast_width.py`	Tests bitcast width enforcement.
`platform/tests/core/test_arch_mma_ssot.py`	Tests integer WMMA dtype aliases.
`platform/tests/core/test_arch_domain_artifact.py`	Tests ROCm compiler discovery.
`platform/python/rocke/tests/test_ir_validity.py`	Tests compile-and-link verdicts.
`platform/python/rocke/tests/ir_validity.py`	Adds the IR-validity test helper.
`platform/python/rocke/tests/helpers/tiling/test_store_fragment_casts.py`	Tests store conversion rejection.
`platform/python/rocke/tests/helpers/tiling/test_fragment.py`	Tests fragment register invariants.
`platform/python/rocke/tests/helpers/tiling/test_fill_fragment.py`	Uses real IR values in fill tests.
`platform/python/rocke/tests/gpu.py`	Adds architecture-aware GPU testing.
`platform/python/rocke/helpers/tiling/transforms/verb.py`	Updates transform documentation.
`platform/python/rocke/helpers/tiling/mma/driver.py`	Documents fragment and return contracts.
`platform/python/rocke/helpers/tiling/fragments.py`	Enforces fragment type and value invariants.
`platform/python/rocke/helpers/tiling/emit.py`	Validates fragment store conversions early.
`platform/python/rocke/helpers/tiling/docs/tiling_api_surface.md`	Documents mutable fragment values.
`platform/python/rocke/helpers/tiling/docs/tiling_api_contract.md`	Corrects rank-validation ownership.
`platform/python/rocke/core/tf32.py`	Derives TF32 support from catalogs.
`platform/python/rocke/core/storage.py`	Adds register-packing utilities.
`platform/python/rocke/core/lower_llvm.py`	Lowers F64 and typed MMA operands.
`platform/python/rocke/core/lower_hip.py`	Uses catalog-driven TF32 validation.
`platform/python/rocke/core/isa/backend.py`	Types WMMA operands at intrinsic boundaries.
`platform/python/rocke/core/ir.py`	Adds types, width checks, and F64 operations.
`platform/python/rocke/core/dtypes.py`	Defines aliases and widths for new types.
`platform/python/rocke/core/arch/data/arch_specs.json`	Canonicalizes integer WMMA operand types.
`platform/python/rocke/core/__init__.py`	Exports new core types.
`platform/python/rocke/__init__.py`	Exports new public types.
`platform/dsl_docs/reference/dtype_support.md`	Documents dtype coverage and engine gaps.
`platform/dsl_docs/ir_lowering/backend_details.md`	Updates LLVM type documentation.
`platform/dsl_docs/development/extending.md`	Adds the dtype-extension procedure.
`platform/dsl_docs/architecture/dtype_storage_packing.md`	Documents register packing.

🧠 Review effort: Balanced

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

tenpercent · 2026-10-10T04:58:20Z

note: the description looks cut in a few places. Please also see #13141 for clash in f64 base support

Copilot

🔵 Needs a closer look

Broad cross-engine changes need human sign-off, resolution of correctness findings, and refreshed validation after the pending rebase.

3 open findings

5 resolved since last review

🧠 Review effort: Balanced

Copilot

🟡 Changes recommended

Missing newer-flavor artifacts and host-dependent test failures leave validation gates broken on supported configurations.

4 open findings

2 resolved since last review

🧠 Review effort: Balanced

Copilot

🔵 Needs a closer look

Shared type contracts and cross-architecture compiler changes require human confirmation of the reported native, parity, and GPU validation.

No open issues 🎉

4 resolved since last review

_{Review effort: Balanced}

tenpercent

Reviewed head a0096d5e819a6bbb70f8786dbfe99319cb50c562. Approving the current scope, with the following non-blocking follow-up suggestions:

Logical i4/u4 tensors and fragments. The documented packed-only scope is acceptable for this PR. A follow-up should separate logical element dtype and coordinates from byte/word transport, then define packed addressing, alignment and partial-store ownership. Packed storage does not preclude logical i4 tensors: ROCm clang 20 compiled and linked a <8 x i4> global -> LDS -> global copy for gfx942, emitting dword global accesses and 32-bit LDS reads/writes. Please also correct the overly broad statement in platform/python/rocke/helpers/tiling/fragments.py that LLVM cannot pack sub-byte formats.
Low severity: validate uitofp_f32 at lowering boundaries. Normal builder-produced kernels are unaffected, but hand-built or deserialized IR can bypass the unsigned-source guard. Retyping a serialized u8 input to i8 passes parsing and verification, and both engines' LLVM/HIP lowerers accept it: bits 0xff mean 255.0 under LLVM uitofp i8, but -1.0 under HIP's signed int8_t cast. Add source-type validation in both engines' lowerers and a serialized-input refusal test. This can be handled in a narrow follow-up.
Low severity: installed native-builder parity wiring. The installed rocke_pytest environment in platform/CMakeLists.txt omits ROCKE_IR_TYPES_TEST, so its native-builder comparison cases skip. Supply the installed fixture path and verify the generated environment in provider and standalone layouts. The comparisons passed locally when the fixture paths were supplied.

Validation at this head: 588 tests passed in strict Python/C++ mode (ROCKE_BACKEND=both ROCKE_CPP_STRICT=1), with 16 skips; six focused C++ test executables passed; 117 catalog tests passed. GPU numerics, gfx1250 compilation and full domain-artifact regeneration were not verified in this review. These suggestions do not block approval of the documented scope.

CongMa13

No new issues found. I agree with the following existing comments:

Mirror canonical operand rows in C++ tables — copilot-pull-request-reviewer, #13434 (comment)
Add new public scalar types to the C++ registry — copilot-pull-request-reviewer, #13434 (comment)
Mirror operand bitcast handling in the C++ lowerer — copilot-pull-request-reviewer, #13434 (comment)
Align f64 allocations to 8 bytes — copilot-pull-request-reviewer, #13434 (comment)
Preserve conversion edge cases in GPU assertions — copilot-pull-request-reviewer, #13434 (comment)
Use unsigned arithmetic semantics for U8 — copilot-pull-request-reviewer, #13434 (comment)
Preserve opcode values or change the ABI — copilot-pull-request-reviewer, #13434 (comment)
Use the selected device architecture in GPU tests — copilot-pull-request-reviewer, #13434 (comment)
Add U8 min/max declarations to each LLVM architecture-domain artifact — copilot-pull-request-reviewer, #13434 (comment)
Skip GPU tests when HIP or a device is unavailable — copilot-pull-request-reviewer, #13434 (comment)
Skip the ROCm compiler lookup test when installed ROCm has no clang — copilot-pull-request-reviewer, #13434 (comment)

CongMa13

No new P0 or P1 issues were identified during this review. The changes are approved.

…y compiler check_ir_validity._llvm_tool looked for clang beside the comgr rocke had loaded, then in $ROCM_PATH/$ROCM_HOME, then on PATH. When the loaded comgr is one bundled with another package (the ROCm torch wheel ships its own), no LLVM sits beside it and the lookup fell through to PATH, typically a distro clang that cannot link AMDGPU code objects. The lookup now also searches the real ROCm installs rocke's runtime already discovers before falling back to PATH. Co-Authored-By: Claude <noreply@anthropic.com>

rocke.tests.ir_validity lets a test ask whether lowered LLVM IR compiles and links for a target, using the same compiler lookup and compile command as platform/tools/check_ir_validity.py. It reports accepted, rejected (with clang's first error line) or cannot_judge (no tools, no clang, a clang that does not know the target, or a clang that cannot run its linker); assert_target_accepts turns cannot_judge into a named skip. Co-Authored-By: Claude <noreply@anthropic.com>

cvt_f32_to_f64 (exact, fpext) and cvt_f64_to_f32 (round to nearest even, fptrunc) are builder methods on their own opcodes, arith.cvt_f32_to_f64 and arith.cvt_f64_to_f32, in both engines: the Python and C++ builders, the LLVM lowerings and the HIP lowerings. Each lowers to one instruction or one C cast. cast_to_f32 and cast_f32_to remain the f16/bf16 <-> f32 hub and keep rejecting f64, so no code built on them converts a double by accident; a conversion given the wrong source type fails with the same message in both builders and emits no op. zero_vec accepts F64 and I16 in both builders, so fill_fragment can zero an f64 fragment, and smem_load_vN accepts i16 in both builders, as smem_store_vN already did. rocke.tests.gpu runs a numeric test on the current GPU, the one rocke's runtime launches on, which hip_module now reports through a public current_device(), only when its arch is one the test names as supporting its data types, lowering and compiling with the active engine; on such an arch a lowering error or compiler rejection fails the test. On the GPU, copy, neg, abs, add, sub and mul of f64 and both conversions are checked bit-exact over ordinary values and the IEEE edge values (out-of-f32-range, subnormal, signed zeros, infinities, NaN), except where IEEE 754 leaves the bits open. The C++ builder side is covered by the rocke_f64_lowering ctest. The two opcodes are appended to rocke_opcode_t, so no existing opcode value changes, and the rocke_opcode_values ctest pins the value, name and purity of every opcode that existed before them. Co-Authored-By: Claude <noreply@anthropic.com>

…nds i8 and i4 I4 is a logical 4-bit integer type, exported next to I8. It is packed-only: its values only ever live packed in 32-bit words, so it has no scalar or vector form and no LLVM or HIP rendering. The 8- and 4-bit integer dtypes are i8 and i4. RDNA integer WMMA instructions say iu8/iu4 in their names, but signedness is an instruction operand, not a property of the data, so iu8 and iu4 are no longer dtypes: their dtype entries are removed and no alias replaces them. The gfx1151 and gfx11-generic integer WMMA rows now record i8 and i4 operands; their op ids are unchanged. A query for i8 or i4 on those targets finds those instructions; a query by iu8 or iu4 finds nothing. Both engines make both changes: the C++ engine gains rocke_i4(), its dtype table drops iu8/iu4, and its embedded catalog rows match the JSON ones. A C++ test pins the catalog lookups: i8/i4 find the ops, iu8/iu4 do not. The HIP lowerings refuse I4 the way the LLVM lowerings do. Every HIP type lookup goes through one function per engine (_hip_scalar_type in lower_hip.py, rocke_h_require_hip_scalar in the C++ lowerer), which refuses a type with no HIP spelling with NotImplementedError / ROCKE_ERR_NOTIMPL "no HIP mapping for type <name>", replacing a bare KeyError in Python and four differently worded C++ messages. The C++ HIP lowerer now records the reason on the builder's error channel, as its header documents. The IR builder of both engines enforces the packed-only rule: any op result, kernel parameter or LDS tile of i4 elements is refused with the TypeError "<op>: i4 values travel packed in 32-bit words and have no scalar or vector form -- keep them in i32 storage and unpack explicitly". The check lives in one function per engine, _refuse_packed_only_value in core/ir.py and rocke_i_refuse_packed_only_value in the C++ builder, called for every op result and every parameter; zero_vec calls it before its own element list. The lowerers' scalar and LDS type tables carry a comment naming it as the reason they have no i4 entry, and their own refusal stays as a backstop. rocke.tests.engines lowers a kernel with each engine and requires the same IR, or the same error, whatever backend the run selects, for the LLVM and the HIP lowering; the rocke_engine binding gains lower_serialized_ir_hip for the second. test_packed_only.py pins the Python builder's refusal text for each builder that could make an i4 value, and uses rocke.tests.engines to check that both engines refuse an i4 parameter, assembled by hand past the builder, in both lowerings. dtype_support.md names the integer dtypes, describes the packed-only rule, and rescores i4. Co-Authored-By: Claude <noreply@anthropic.com>

U8 and U4 name unsigned 8- and 4-bit integer data (spellings u8/uint8 and u4/uint4), exported next to I8 and I4. They let a kernel say its data is unsigned so that matrix instructions with a per-operand signedness flag can be selected for it. U8 renders as LLVM i8, since LLVM integers carry no sign, and as HIP uint8_t. U4, like I4, is packed-only: the IR builder of both engines refuses a u4 value or LDS tile with the I4 refusal's text (_refuse_packed_only_value, rocke_i_refuse_packed_only_value), and it has no LLVM or HIP rendering, so no u4 value reaches an arithmetic op or a compare. U8 is an arithmetic type. One predicate per engine says whether a type is unsigned (_is_unsigned in core/ir.py, rocke_type_is_unsigned in the C++ engine), and every op whose result depends on signedness asks it which form to emit for u8: div and mod lower to udiv and urem, the ordered compares of arith.cmp and vector.cmp to icmp ult/ule/ugt/uge, and add, sub and mul drop the nsw flag, since unsigned values wrap. zext, lshr and the sign-agnostic ops lower as for i8. Three ops are new: umax and umin (llvm.umax.i8 / llvm.umin.i8, measured into the llvm20, llvm22 and llvm23 arch-domain artifacts with those toolchains), uitofp_f32, and cvt_f32_to_u8_sat, which rounds to nearest even and clamps to [0, 255] in f32 before fptoui, so a NaN gives 0. smax, smin, sext and their vector forms are signed by definition. Both builders refuse a u8 operand when one of them is built, vector_binary("smax" | "smin") included, naming the unsigned op to use where there is one. The LLVM lowerings of both engines, and their HIP lowerings of the scalar ops (HIP has no u8 vector type), refuse it again with the same text under the op's dotted name (_refuse_unsigned_operands, rocke_i_refuse_unsigned_operands), so IR that did not come through a builder cannot run a signed op on unsigned data. The new opcodes are appended to rocke_opcode_t, so no existing opcode value changes. In HIP every u8 value is a uint8_t, so C's own operators are the unsigned ones, and umax / umin share smax / smin's ternary without their refusal; zext casts a u8 source through uint8_t. dtype_support.md describes what u8 supports. A u8 LDS element takes one byte, like i8: the lowerers' element-size and alignment tables list u8 next to i8, so a u8 tile and its scalar stores lower exactly as the i8 ones do. smem_load_vN and smem_store_vN admit one shared list of vector element types (_SMEM_VECTOR_ELEMENTS in core/ir.py, ROCKE_SMEM_VECTOR_ELEMENTS in ir_tile.cpp), refused with one message generated from it. smem_store_vN had no element check and sized an unlisted element as two bytes; it now refuses a u8 vector as smem_load_vN and the global vector builders do. Both engines make every change. Tests lower a U8 copy and a U8 LDS store with both engines, to LLVM IR (the LDS store matches the I8 kernel) and to HIP source (compiled with hipcc when present). test_u8_arith.py lowers every u8 op with both engines and checks the unsigned LLVM forms and the HIP types, and on a gfx90a GPU checks division, remainder, the compares, umax/umin, lshr, wrapping add/sub/mul, uitofp, zext and the saturating conversion bit-exact against numpy uint8, on values above 127. Tests also check that vector_binary refuses smax and smin on u8, and that both engines refuse each signed op on u8 in deserialized IR, in the LLVM lowering and, for the scalar ops, the HIP lowering. test_packed_only.py checks the builders' refusal of u4 values, and that both engines refuse a u4 parameter in both lowerings. A new ctest, rocke_ir_types, builds LDS vector accesses, i4 and u4 values and the u8 ops and refusals with the C++ builders; tests compare its answers with the Python builders, and tests/run_all.py passes it in as ROCKE_IR_TYPES_TEST. Co-Authored-By: Claude <noreply@anthropic.com>

A bitcast keeps every bit of its operand, so a source and target of different widths is always a mistake; until now it was caught only by clang ("invalid cast opcode"). type_bits(t) in core/ir.py is the one place that answers how many bits a value of an IR type occupies: a scalar's width comes from its format, a vector's is its count times its element width, and a pointer's comes from its address space (32 bits for LDS, 64 otherwise). An LDS allocation is memory, not a value, and raises TypeError. IRBuilder.bitcast and IRBuilder.vec_bitcast raise ValueError naming both types and widths when the widths differ, and TypeError for a pointer operand or target, which LLVM does not bitcast. Equal-width casts lower exactly as before, so no existing kernel's IR changes. The C++ builders refuse the same casts with the same text: rocke_b_bitcast and rocke_b_vec_bitcast call rocke_i_check_bitcast_width before building the op. rocke_i_type_bits is the C++ twin of type_bits, and rocke_require_dtype_info (rocke/dtypes_internal.h) is the raising twin of rocke_dtype_info, so "unknown dtype '<name>'" has one home in each engine. The rocke_ir_types test binary gains a --bitcast mode; the tests require its answers to equal the Python builder's, and lower the equal-width casts with both engines. Co-Authored-By: Claude <noreply@anthropic.com>

register_packing(dtype, count) in core/storage.py returns the dense FragmentPacking for count elements of a sub-byte dtype (fp4, fp6, i4, u4) held in whole 32-bit registers: elements pack low-bit-first with no padding, so 32 fp4 values take 4 registers and 32 fp6 values take 6. A payload that leaves a register partly filled raises ValueError, a predicate raises, and a dtype of 8 or more bits returns None: it has an LLVM element type, and LLVM packs its vectors into registers itself. packed_word_range(packing, start, length) returns the range of registers holding elements [start, start + length), and raises when the range does not start and end on a register boundary (naming the element step that does) or runs past the fragment. Both reject a non-integer count, start or length with TypeError. rocke_register_packing and rocke_packed_word_range in rocke/storage.h are the C++ twins. They keep the header's promise that no exception crosses the C API: they return false and write Python's refusal text into a caller-owned buffer. The rocke_storage test binary answers both questions on the command line, and test_storage.py requires its answers and texts to equal core/storage.py's. Both are pure arithmetic with no caller yet. dtype_storage_packing.md describes them. Co-Authored-By: Claude <noreply@anthropic.com>

…dtype test_dtypes.py gains _IR_SCALAR_BITS, the width of every IR scalar type written out by hand. The format table in core/dtypes.py and type_bits are both checked against it, and a completeness test fails when an IR scalar type is added without its width. The scalar width check in test_bitcast_width.py moves here. The same map checks the C++ engine: rocke_ir_types --scalar-bits prints every IR scalar type with the width rocke_i_type_bits gives it. Its name list holds one name per scalar kind, which a static_assert and a self-check enforce, so a new kind cannot go unchecked. extending.md section 5.4 becomes the one recipe for adding a dtype: format table, IR type, reference width, register form, memory, conversions, MMA catalog, manifest, the C++ engine mirror, docs and tests. backend_details.md points to it instead of keeping its own checklist, and its list of scalar LLVM mappings gains i16, f64 and u8. Co-Authored-By: Claude <noreply@anthropic.com>

Four MMA lowerings printed a fixed register type around whatever value they received: the two f32 MFMAs print `float`, the gfx11 integer WMMA prints `<N x i32>` and the gfx1250 fp8 WMMA prints `<8 x i32>`. A value of the instruction's own format in another shape, such as a one-element `<1 x float>` vector or packed `<16 x i8>` / `<32 x i8>` bytes, produced IR that LLVM rejects. The HIP lowering passed the same values to the builtins unchanged: clang rejects the one-element vector for an f32 MFMA, and accepts the WMMA bytes only through an implicit vector conversion. core/mma_operand.py holds the rule, with no output syntax. mma_operand_needs_reinterpret(op_id, role, operand_type, register_type, arch) returns False for an operand of the register type, without consulting the MMA catalog, and True when one reinterpretation of the same bits gives the register type: the operand holds the register's element type or the element type the target's MMA catalog records for that operand, and has the register's bit width. Any other number format, and any other bit width, raises ValueError naming the operand, the instruction and the register type. mma_operand_value(op, role) picks operand a or b. Each lowerer only renders the answer. In the LLVM lowering, _Lowerer._mma_operand(op, role, register_type) emits one `bitcast` through _Lowerer._reinterpret; _mma_op_id reads the op_id from an MMA op's `tile.<op_id>` name. In the HIP lowering, _Lowerer._mma_operand(op, op_id, role, register_type) writes `__builtin_bit_cast(<register>, <operand>)`. The four lowerings pass their operands through them in both lowerings. Operands that already have the register type lower exactly as before, so no existing kernel's IR or HIP source changes. Tests cover each lowering with different A and B values, the number-format and width errors for both operands with the same text from the LLVM and HIP lowerings, pin the gfx11 integer WMMA signedness and clamp flags, and compile the HIP source with hipcc when it is present. The C++ engine applies the same rule with the same fresh names and error text. rocke_mma_operand_needs_reinterpret and rocke_mma_operand_value (cpp/core/mma_operand.cpp, rocke/mma_operand_internal.h) are the twins of the rule module; mma_operand in lower_llvm/mma.cpp renders through rocke_ll_reinterpret, and the HIP lowering of the two f32 MFMAs renders through h_mma_operand in lower_hip_mma.cpp. rocke_require_ir_type is the raising twin of rocke_dtype_to_ir_type, with dtype_to_ir_type's texts, and rocke_i_vector_type_in_arena builds the register type in the lowerer's arena. Every LLVM test lowers with both engines and requires the same IR, or the same error; the f32 MFMA HIP tests do the same for the HIP source. Co-Authored-By: Claude <noreply@anthropic.com>

IRBuilder.mma chose the result element as i32 for integer accumulators and f32 for everything else, so an instruction with any other accumulator type would have been given the wrong result type. _mma_c_dtype(op_id) replaces _mma_c_is_int and returns the accumulator dtype the arch catalog records for an op_id (fp32 for op_ids the catalog does not list). IRBuilder.mma types its result as dtype_to_ir_type of that dtype, or of the MmaOp's own c_dtype when it is given one. Every catalog instruction today accumulates in fp32 or i32, so no kernel's IR changes. Tests check that every catalog accumulator has an IR type, that an op_id without a catalog row accumulates in fp32, and that IRBuilder.mma results take the accumulator dtype for an fp32 and an i32 instruction. The IRBuilder.mma docstring describes the result type. The C++ rocke_b_mma types its result the same way: rocke_mma_c_dtype replaces rocke_mma_is_int_acc, and rocke_require_ir_type resolves the dtype, so a refusal has core/ir.py's text. It now checks the operands before the op_id, in IRBuilder.mma's order (scaled MFMA, then XF32), so both builders refuse the same call with the same text. The scaled MFMA check gives one refusal for a wrong operand count or type, as mfma_scale.py does, instead of a separate "requires five operands". The rocke_ir_types test gains --mma <op_id>, which builds one MMA and prints its result type or the refusal; test_mma_frag_tables.py compares it with the Python builder for every op_id in the catalog and the fragment table, and checks that both refuse an unknown op_id. Its self-check pins that a refusal builds no op. The rocke_mma_frag_ssot test adds an op_id with no catalog row and checks every catalog accumulator. Co-Authored-By: Claude <noreply@anthropic.com>

The LLVM and HIP lowerings refused a TF32 MMA on any target other than gfx942 by comparing the target name, with the message "XF32 MMA requires gfx942". The arch catalog already records which targets list each XF32 instruction. tf32_mma_target_error(op_id, arch) in core/tf32.py returns the error for a TF32 op_id that the target's MMA catalog does not list: "no MMA instruction found for '<op_id>' on '<target>' -- the target's MMA catalog does not list it; it is listed for <targets>". The three call sites (the neutral MMA and XF32 LLVM lowerings, and the HIP MMA lowering) raise it as before, as a ValueError; the lowerers import core/tf32.py at module level, as it imports nothing that imports them. gfx942 lowers exactly as before, so no kernel's IR changes, and a target that gains the instruction in its catalog needs no code change. The C++ LLVM and HIP lowerers decide the same way through rocke_tf32_mma_target_error, with the same text. It is declared in tf32_internal.h and defined in the new cpp/core/tf32.cpp, so its catalog lookup is compiled once instead of inlined into every file that includes the header. The tests require the exact message from both LLVM lowerers and from the C++ HIP lowerer. dtype_support.md describes the rule. Co-Authored-By: Claude <noreply@anthropic.com>

assistant-librarian · 2026-10-22T19:20:55Z

TheRock Submodule Bump Activity

Newest first

2026-10-22 — Included in TheRock via ROCm/TheRock#9426.

ISSUE ID : [AICK-2640](https://amd-hub.atlassian.net/browse/AICK-2640) Related base f64 support: [AICK-1713](https://amd-hub.atlassian.net/browse/AICK-1713). ## Motivation #13141 added the f64 type, memory and lowering, and #13434 made MMA result and operand typing catalog-driven. This PR adds the two gfx942 f64 MFMA atoms, `v_mfma_f64_16x16x4_f64` and `v_mfma_f64_4x4x4_4b_f64`, so kernels can issue f64 matrix instructions through `IRBuilder.mma`. It is the step before f64 UniversalGemm. ## Technical Details All changes are made in both the Python and C++ engines. - **Catalog:** gfx942 rows `mfma_f64_16x16x4_f64` (16x16x4) and `mfma_f64_4x4x4_f64` (four 4x4x4 blocks) in `arch_specs.json` and `k_mma_gfx942`. No other arch lists them, and no GEMM family selects them yet: every family asks the catalog for an fp32 accumulator, and UniversalGemm still requires fp32 accumulation. - **Lane layouts:** - 16x16x4: A and B use the same lane maps as fp32 16x16x4. The shared maps are renamed to `_mfma_a_16x16x4` / `_mfma_b_16x16x4`. - The 16x16x4 accumulator has its own layout, `_mfma_acc_16x16x4_f64`: slot `i` of lane `l` holds row `4*i + l/16`, column `l%16`. A lane's slots are 4 rows apart, unlike the contiguous rows of `_mfma_acc_16x16`. - 4x4x4: the four blocks are stacked along rows, so A, B and C/D are 16x4. A reuses `_mfma_a_16x16x4`; B and C/D have their own maps (`_mfma_b_4x4x4_f64`, `_mfma_acc_4x4x4_f64`). - **Result and operand types:** results follow the catalog accumulator dtype (#13434), so the f64 atoms give `<4 x f64>` and `<1 x f64>`. A and B pass through the `mma_operand` rule like the f32 scalar atoms: an `f64` value passes as is, a one-element `vec<f64x1>` is reinterpreted with one bitcast, and any other number format is rejected. - **LLVM lowering:** declarations for `llvm.amdgcn.mfma.f64.16x16x4f64` and `llvm.amdgcn.mfma.f64.4x4x4f64`. The 4x4x4 intrinsic takes and returns a scalar `double`, so the lowering extracts the scalar from the `<1 x double>` accumulator and re-inserts the result. - **HIP lowering:** `__builtin_amdgcn_mfma_f64_16x16x4f64` / `__builtin_amdgcn_mfma_f64_4x4x4f64`, plus an `f64x4` vector typedef. - **Probe instance:** `instances/gfx942/f64_mma_probe` (`F64MmaProbeSpec(m, k)`, m = 16 or 4): one wave computes D = A·B + C, loading and storing through the catalog lane maps. The C++ mirror is exposed through `rocke_engine.f64_mma_probe_serialize_ir`. - **Intrinsic arch-domain artifacts:** `intrinsic_arch_domain.llvm{20,22,23}.json` regenerated with `tools/gen_arch_domain.py` on each flavor's toolchain. They add rows for the two new intrinsics; no existing row changes. - **Docs:** `mfma_atom_catalog.md` (atoms and f64 output layouts), `mma_metadata.md` (f64 operand and result types) and `dtype_support.md`. ## Test Plan - Byte-identity gate on llvm20, llvm22 and llvm23, with a new parity family `f64_mma` (m = 16 and 4; k = 4, 8 and 64). - `tests/core/test_f64_mma.py`: catalog rows, lane-map formulas, result types, and Python/C++ IR equality. - `tests/numeric/test_gfx942_f64_mma_numeric.py`: both atoms through both engines on gfx942, against a float64 reference. - f64 rows in `test_mma_frag_tables.py` and the `rocke_mma_frag_ssot` ctest. - `test_arch_domain_artifact.py` on all three flavors; ctest and the full pytest suite compared against `develop`. ## Test Result - Byte-identity: GREEN on llvm20, llvm22 and llvm23, including `f64_mma`. - f64 MMA tests pass; the GPU results match the reference exactly on integer inputs. - Arch-domain artifact test and ctest pass; no new pytest failures. Next step: f64 UniversalGemm on gfx942. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. [AICK-2640]: https://amd-hub.atlassian.net/browse/AICK-2640?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ [AICK-1713]: https://amd-hub.atlassian.net/browse/AICK-1713?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ --------- Co-authored-by: Maksim (Max) Podkorytov <Maksim.Podkorytov@amd.com>

Copilot AI balanced review requested due to automatic review settings October 9, 2026 23:10

cgmillette requested review from a team as code owners October 9, 2026 23:10

github-actions Bot added documentation project: hip-kernel-provider rocke: platform Touches rocKE platform labels Oct 9, 2026

Copilot started reviewing on behalf of cgmillette October 9, 2026 23:11

View session

assistant-librarian Bot added the organization: ROCm label Oct 9, 2026

Copilot AI reviewed Oct 9, 2026

View reviewed changes

cgmillette mentioned this pull request Oct 13, 2026

feat(rocke): add FP64 support for elementwise and reduce #13141

Merged

1 task

cgmillette changed the title ~~feat(rocke): add f64, i4 and unsigned types, bit-width and packing contracts, and MMA operand typing~~ feat(rocke): add i4 and unsigned types, bit-width and packing contracts, and MMA operand typing Oct 13, 2026

cgmillette changed the title ~~feat(rocke): add i4 and unsigned types, bit-width and packing contracts, and MMA operand typing~~ feat(rocke): add i4 and unsigned types, explicit f64 conversions, bit-width and packing contracts, and MMA operand typing Oct 13, 2026

cgmillette force-pushed the users/cgmillette/rocke/core-dtype-foundation branch from 1b714bd to 49535fa Compare October 15, 2026 14:08

cgmillette requested review from CongMa13 and tenpercent and a balanced review from Copilot October 15, 2026 14:14

Copilot started reviewing on behalf of cgmillette October 15, 2026 14:15

View session