Skip to content

fix(rocke): make both engines emit the same HIP source - #14845

Open
cgmillette wants to merge 12 commits into
developfrom
users/cgmillette/rocke/hip-text-identity
Open

cgmillette wants to merge 12 commits into
developfrom
users/cgmillette/rocke/hip-text-identity

Conversation

@cgmillette

@cgmillette cgmillette commented Nov 1, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Both rocKE engines lower a kernel to HIP source as well as to LLVM IR. The LLVM-IR route is held byte-identical by the byte-identity gate; the HIP route had no gate and had drifted. Lowering every parity kernel through both HIP lowerers on develop:

Outcome Platform parity (803 kernels) Library parity (963 kernels)
Same text 377 45
Both lower, text differs 147 144
Python lowers, C++ refuses 241 352
Both refuse 38 422

The differences, by class (a kernel can be in more than one):

Class Platform Library Example (Python / C++) Effect
float literal, whole number 134 132 0.0f / 0f C++ output does not compile
float literal, precision 31 76 -1.4426950408889634f / -1.4427f C++ output changes the constant
double literal, whole number 7 2 0.0 / 0 same value, different text
s_waitcnt on a target with split wait counters 1 0 49279 / 64519 neither compiles: clang cannot select it on that target
C++ has no HIP lowering 241 352 f16/bf16/fp8/bf8 MFMA atoms, gfx11, gfx12 and newer WMMA atoms, tile.inline_asm (now lowered by develop, #12167) C++ refuses
whitespace / ordering 0 0

Technical Details

  • Float literals (C++ changes). The C++ HIP lowering printed float constants with %g. It now prints Python's repr(float), the shortest decimal that reads back as the same double. One formatter, rocke_append_python_float_repr (cpp/include/rocke/float_format.h, cpp/support/float_format.cpp), now serves the C++ IR printer, the IR serializer and the HIP lowerer, and replaces the two copies the printer and serializer kept. It takes the shortest round-trip digits from std::to_chars, so it matches repr exactly (the old copies printed 2**-24 with 17 digits) and does not depend on the C locale. Each engine's HIP lowering adds one thin adapter on top: an f suffix for float literals and the NAN / INFINITY macros for the values repr spells in a way C++ does not accept. Python's spelling is the one kept: it is valid C++ and exact.
  • s_waitcnt on a target with split wait counters (both change). A newer target has no monolithic s_waitcnt. Both HIP lowerings now ask the target's ISA backend (the fact the LLVM route already uses) and refuse with the same message. The C++ HIP arch seam records this as ROCKE_HIP_WAITCNT_NONE. Every other target keeps its encoding. Nine kernels on that target (1 platform, 8 library) that Python used to lower to uncompilable source are now refused by both.
  • Missing C++ lowerings (C++ changes). The f16/bf16/fp8/bf8 MFMA atoms and the WMMA atoms are ported from their Python handlers: one table of op_id-to-builtin names and one small emitter per operand form. Accumulator widths and packed operand word counts come from the arch MMA catalog. The gfx11 integer WMMA atoms lower under every signedness row of the catalog, and their builtin's signedness bools come from the row through rocke_matrix_operand_is_signed, which the LLVM matrix-call lowering now uses as well (the twin of Python's Signedness.flag). The WMMA gate now admits an atom by the target's MMA catalog, as Python's does; the unused has_wmma and family fields of rocke_hip_arch_t are removed, which changes its size, so ROCKE_ABI_VERSION and BINARY_ABI go to 9 (develop + 1).
  • MMA arity (both change). A kernel read from serialized IR is not verified. Every path of both HIP lowerings that lowers an atom from A, B and C now checks first that the op has three operands and one result and refuses it otherwise with the same ValueError (tile.mma <op_id> expects 3 operands and 1 result); the C++ fp8 and integer WMMA emitters, the f32/f64 MFMA controls and the TF32 atoms used to index the arrays unchecked. The Python scaled WMMA lowering counts its five operands with the C++ message. An op_id neither engine lowers is still refused as unknown.
  • Inline asm (both change). Develop's lowering of tile.inline_asm (feat(rocke): add NVFP4 GEMM with tensor scales and focused example #12167) is kept as it is. Both engines now read its template and constraints attrs through one helper each (_inline_asm_string_attr, h_inline_asm_string_attr): a missing one is refused with KeyError tile.inline_asm: missing '<key>' attr and a non-string one with TypeError HIP inline asm <key> must be a string, by both engines (Python used to raise a bare KeyError, or a TypeError from re.sub).
  • LDS access before its allocation (Python changes). The Python HIP lowering raised RuntimeError for an op on an LDS allocation not lowered yet; C++ refuses with the same message as ValueError (there is no RuntimeError status). Python now raises ValueError.
  • Gate. platform/tests/instances/differential/test_hip_source_identity.py runs every platform and library parity emitter (through the differential gate's batch worker) and lowers each kernel through both HIP lowerers: the text must match, or both must refuse with the same error type and message. The serialized IR does not carry the target, so the emitters gain a target-ir mode (target <gfx> line, then the serialization). The test fails if an emitter does not accept the mode or builds no kernel (conv_abi is listed as building none), and fails any config that exits non-zero after printing, exits 0 without printing, or fails with an exception other than a refusal type (ValueError, TypeError, KeyError, NotImplementedError); a config that fails with no output counts as the builder refusing it only if the ir mode refuses it with the same error.
  • Gate in installed CI. CI runs installed CTest, without the rocke_engine binding. tests/core/ir_lower_cli.cpp gains --hip-batch, which lowers a list of kernels to HIP in one run, and the test starts each family's emitter run ahead of the test that reads it; it is built and installed as rocke_ir_lower_cli (not a test itself), and rocke.tests.engines lowers through it when the binding is absent (ROCKE_IR_LOWER_CLI). The rocke_hip_source_identity CTest runs the gate there (120 s per family, 300 s in all) and is in the quick and ffm-quick tiers. It is not registered on Windows, whose CI shards run close to their time limit, so CI does not compare the MSVC-built engine's HIP output with Python's (it matched on every parity kernel on the Windows lanes when the gate was added). rocke_pytest does not collect it.
  • Serialized-IR parse speed (C++ changes). The C++ parser of serialized IR looked up each operand name by scanning every name defined before it, which is quadratic in the kernel's size and was most of the gate's C++ time. It now uses a hash table; parsed kernels are unchanged, and test_serialized_ir_name_scopes.py holds both engines to the same output for a kernel whose sibling loops reuse their names.
  • Focused tests. tests/core/test_hip_lowering_parity.py lowers every native MFMA and WMMA atom of each target with its own catalog (operands built from zero words, so the fp6/fp4 and unsigned formats are included), and the 4x4x4 atoms through both engines (same source or same refusal), refuses every atom without its accumulator or result and every scaled atom without its scales alike in both, checks the type and message of the s_waitcnt refusal on the target with split wait counters and of the inline-asm template / constraints refusals, and holds rocke_ir_lower_cli to the binding. The HIP cases of test_integer_wmma_signedness.py and test_mma_operand_register_type.py now run both engines. The both-engine HIP helpers live in rocke.tests.engines (hip_outcomes_of_both_engines, HipRefusal). tests/core/test_float_format.py and the rocke_float_format ctest check the formatter against repr (a table of hard cases, about 150k sampled values, and every HIP literal spelling).

LLVM-IR output is unchanged.

Test Plan

  • Build engine and binding; ctest.
  • Byte identity at llvm20, llvm22, llvm23, llvm24; golden check (ll at all four flavors, canonical IR).
  • New HIP identity test on this branch, and against the develop C++ engine (must fail).
  • hipcc --genco of every kernel whose C++ HIP text changed (up to two per family, arch and kind): gfx942, gfx950 and gfx90a with ROCm 7.1; gfx1151, gfx1201 and gfx11-generic in the CI image.
  • Full platform pytest suite, default and ROCKE_BACKEND=both strict, compared with develop by failure name; library tests that run parity emitters.

Test Result

Head after the rebase on develop (#12167, #14835 and the rest):

  • Build OK; ctest: 0 of 48 failed (the arch-domain sweep skips, as on develop).
  • Byte identity at llvm20, llvm22, llvm23 and llvm24: GREEN. LLVM-IR output is unchanged.
  • HIP identity test with the binding, ROCKE_BACKEND=both strict: 99 families pass. From an installed tree through rocke_ir_lower_cli (ctest -R rocke_hip_source_identity): 100 passed in 47 s, slowest family about 6 s.
  • The PR's tests and the HIP, MMA, inline-asm and differential-harness tests, ROCKE_BACKEND=both strict: pass.
  • Platform suite, ROCKE_BACKEND=both strict: 10853 passed, 1 failed (test_the_cpp_fallback_does_not_warn_twice, which failed on develop in this mode when last compared).
  • pre-commit over the PR range: clean.

Earlier revisions (before #12167 landed; they included this PR's own inline-asm lowering):

  • Build OK; ctest 36/36 (the new rocke_float_format included).
  • Byte identity GREEN at llvm20, llvm22, llvm23 and llvm24. Golden OK (ll x4, canonical IR); no golden file changed.
  • HIP identity test: 95/95 families pass on the current head. Per-kernel counts from the first revision: platform 764 same text, 39 both refuse. Library: 533 same text, 430 both refuse. Against the develop C++ engine it fails on the literal, waitcnt and missing-lowering classes.
  • hipcc (176 kernels, measured on the first revision of this PR): gfx942 14/16, gfx950 105/108, gfx90a 1/1, gfx1151 13/18, gfx1201 15/16, gfx11-generic 2/2 compile, including the fp8 MFMA, WMMA and inline-asm kernels. Every failure is text identical to what develop's Python engine already emits for that kernel (HIP-route gaps outside this change: auto loop carries, atomicAdd overloads, a bf16/fp16 assignment, a 2-byte float atomic); one large fp8 MoE kernel did not finish compiling in 15 minutes and is likewise develop's Python text. On develop the C++ text of the literal class does not compile (invalid digit 'f' in octal constant).
  • Platform suite on the current head (rebased on develop): default 6414 passed, 2 failed; ROCKE_BACKEND=both strict 6413 passed, 3 failed. Every failure also fails on develop: test_binding_compiler_selection, test_installed_pytest_gets_distinct_artifact_environment_variables (develop's installed test environment has a ROCKE_REQUIRE_COMPILER_TOOLS entry that test does not expect) and, with both engines, test_the_cpp_fallback_does_not_warn_twice.
  • Library: test_barrier_fence_parity_emit, test_direct_wgrad_specs, strided-KV emitter golden: pass.

ISSUE ID : AICK-2903

🤖 Generated with Claude Code

@cgmillette
cgmillette requested review from a team as code owners November 1, 2026 22:15
@cgmillette cgmillette changed the title test(rocke): require both engines to emit the same HIP source fix(rocke): make both engines emit the same HIP source Nov 1, 2026
@github-actions github-actions Bot added documentation project: hip-kernel-provider rocke-lib: attention Touches rocKE library attention/FMHA kernels rocke: platform Touches rocKE platform labels Nov 1, 2026
@therock-pr-bot

therock-pr-bot Bot commented Nov 1, 2026 •

Copy link
Copy Markdown

✅ 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

therock-pr-bot Bot commented Nov 1, 2026 •

Copy link
Copy Markdown

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

@cgmillette

Copy link
Copy Markdown
Contributor Author

Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max

PR #14845 — make both engines emit the same HIP source (cgmillette, AICK-2903) — 🟡 First-round: gate holds, one repr gap

Item Status
Author: identity gate, 95 families 🟢 Verified
• 95 ✓ with the engine; each family within CI's per-test timeout (ran)
• 2031 kernels: 0 differ, 0 one-sided refusals (ran)
• local only: not in CI's installed tests, no engine there (read)
Author: develop drift table 🟢 Mostly
• same classes; platform 377 same, 134 whole-number literals (ran)
• totals here 815 + 1216 kernels, body 803 + 963 (ran)
Author: new C++ atoms, WMMA gate, gfx1250 s_waitcnt 🟢 Verified
• 32 kernels × all 8 wired targets: same text or both refuse (ran)
• incl. pairs no parity config has: 32x32x16 fp8/bf8, iu4, MFMA on gfx90a (ran)
• gfx1250: 9 kernels newly refused by both, as stated (ran)
Author: emitted HIP compiles 🟢 Mostly
• all 4223 literal spellings compile; develop's 0f does not (ran)
• 10/17 sampled kernels compile; failures are on Python-identical lines (ran)
Bots / CI 🟢 Green
• gfx94X rocKE shards: develop's golden timeout (14733) on attempt 1; re-run ✓
• 53 ✓; the red Summary is the first run, cancelled by a label re-run
• merge: clean
Our checks 🟢 Clean
• merge vs 92f853c9: clean (ran)
• build + ctest 34 ✓ / 1 skip; byte identity every family ×4 GREEN (ran)
• goldens + differential unchanged; HIP-route tests 3106 ✓ / 72 skip; rules 1–3 clean (ran)
# Sev file:line Finding
N1 🟡 core/ir/print.cpp:77, lower_hip/core.cpp:591 C++ float spelling is not repr for some powers of two
• impact: 46 of 2098 differ, e.g. 2^-24 from the public Philox helper: C++ 5.9604644775390625e-08f, Python 5.960464477539063e-08f; same value, different HIP text (ran)
• fix: shortest round-trip digits (e.g. std::to_chars) in CPython's layout; the IR printer shares it

Utility: 🟢 Positive — engines now emit identical HIP for every parity kernel; gate local-only.


Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted

First-round review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: merge with develop, native build + ctest, byte identity and goldens at four LLVM flavors, the new gate and all HIP-route tests with the engine, a both-engine sweep over every wired target, literal and kernel compiles, CI at head. Flags whether a closer look is needed; not an approval or rejection.

@shumway

shumway commented Nov 2, 2026

Copy link
Copy Markdown
Contributor

This is OK as-is, with the noted N1 formatting finding above.

If you do make changes, I would recommend having agents look at this changes in the context of single-responsibility principle (SRP) and don't-repeat yourself (DRY). This may make the design cleaner and easier to test and maintain. Keeping python and C++ in sync can be subtle, so a good design may help simplify future improvements.

Here's how my agents responded to my questions in SRP and DRY:

What it gets right

  • h_python_float_literal() in platform/cpp/core/lower_hip/core.cpp:551 cleanly separates:
    • Python-compatible numeric spelling.
    • HIP-specific suffixes and special-value spellings.
  • Both rocke_h_f32_literal() and rocke_h_f64_literal() reuse that helper, satisfying DRY within the HIP lowerer.

Problems

  1. SRP/layering: rocke_append_python_float() is exposed from ir_print.h, although it is a general formatting primitive now consumed by the IR printer and HIP lowerer. HIP lowering
    should not depend on the IR-printer interface merely to format a number.

  2. DRY: essentially the same formatter remains duplicated as emit_float_repr() in:

    • platform/cpp/core/ir/print.cpp:63
    • platform/cpp/core/ir/serialize.cpp:137

    serialize.cpp:132-136 explicitly acknowledges the duplication. Exporting only one copy without consolidating the other leaves two implementations that can drift.

  3. Testability: there is no focused unit test for the formatter’s advertised contract. The HIP identity test is corpus-based, so values absent from existing kernels are invisible. The
    2**-24 mismatch demonstrates this gap.

  4. Contract is stronger than implementation: ir_print.h:33-38 promises exact Python repr(float), but the implementation does not always provide it. The %e/strtod approach is also
    locale-sensitive, while Python’s representation is locale-independent.

Better abstraction

Move the primitive to a neutral support module, for example:

  platform/cpp/include/rocke/float_format.h
  platform/cpp/support/float_format.cpp

With responsibilities such as:

  void rocke_append_python_float_repr(rocke_strbuf_t*, double);

Then use that single implementation from:

  • IR printing
  • IR serialization
  • HIP literal formatting

Keep the HIP policy adapter separate:

  rocke_h_f32_literal(...)  // NaN/Inf spelling + "f"
  rocke_h_f64_literal(...)  // NaN/Inf spelling, no suffix

Add a focused table/property test covering:

  • 0.0, -0.0, integral values
  • notation boundaries around 1e-4 and 1e16
  • 2**-24 and adjacent values
  • subnormals and finite extremes
  • NaN and infinities
  • locale independence

So my answer is: the HIP-local helper is reasonable, but the PR does not yet provide the right shared, directly testable SRP/DRY abstraction for float formatting. A small consolidation would materially improve it.

@CongMa13

CongMa13 commented Nov 2, 2026

Copy link
Copy Markdown
Contributor

I reviewed this PR again. The following review threads are still unresolved, or comments have no author reply:

If you have already fixed these, please reply on each comment to say so and resolve the threads.

@cgmillette
cgmillette force-pushed the users/cgmillette/rocke/hip-text-identity branch from 8cc38ce to 83770a6 Compare November 2, 2026 20:26
@cgmillette

Copy link
Copy Markdown
Contributor Author

Thanks @shumway. I took the SRP / DRY suggestion across the whole change, and N1 (2**-24) is fixed. The fixes are folded into the commits they belong to.

Float formatting: one primitive, one adapter per engine

  • rocke_append_python_float_repr (cpp/include/rocke/float_format.h, cpp/support/float_format.cpp) is the only place the C++ engine spells a float repr-style. The IR printer (both attribute paths), the IR serializer and the HIP lowerer call it. The two emit_float_repr copies in print.cpp and serialize.cpp are gone, and ir_print.h no longer exports a formatter. git grep finds one other kind of float printing in cpp/: the %g in conv instance tags. Those mirror Python's f"{x:g}", a different rule. On the Python side, every one of these paths uses repr().
  • The primitive takes the shortest round-trip digits from std::to_chars, with no printf and no strtod probing. It matches repr exactly and does not depend on the C locale. 2**-24 now prints as 5.960464477539063e-08.
  • HIP literal policy is one small adapter per engine: _hip_float_literal in lower_hip.py and h_hip_float_literal in lower_hip/core.cpp. It adds f to a float, nothing to a double, and spells nan and the infinities with the NAN / INFINITY macros. The f32 and f64 entry points are one-line calls into it.

Tests for the formatter

  • The rocke_float_format ctest checks a table whose repr text is known: zeros, integral values, both sides of 1e-4 and 1e16, 2**-24 and its neighbours, subnormals, the finite extremes, nan and the infinities. When a comma-decimal locale is installed, it runs the table again under that locale. CI images have none, so there it checks the C locale only.
  • tests/core/test_float_format.py compares the primitive with repr() over every power of two and of ten, the doubles either side of each, and a fixed random sample (about 150k values). It also checks every float and double literal both HIP lowerers emit.
  • tests/run_all.py and the installed rocke_pytest both pass the binary to it through ROCKE_FLOAT_FORMAT_TEST. The installed-environment test and the run_all fixtures now list it too.
  • I also ran one manual check, which nothing in the branch repeats: the old printf/strtod formatter fails the table on 2**-24, and under a hand-generated de_DE locale it fails 20 cases.

The same treatment for the MMA port and the refusals

  • The C++ MMA atoms are now tables of op_id and builtin name, with one small emitter per operand form. Accumulator widths and packed operand word counts come from the arch MMA catalog (c_frag_len, a_frag_len, b_frag_len) instead of a second width table. The builtin names are clang frontend names that the Python HIP handlers also spell. They stay in the HIP lowering; deriving them from the LLVM lowering's intrinsic table would make the HIP lowerer depend on the LLVM lowerer's internals.
  • rocke_hip_arch_t loses has_wmma and family. Nothing read them after the WMMA gate moved to the catalog. The header comment now describes what the struct actually holds.
  • New tests/core/test_hip_lowering_parity.py. It lowers every native MFMA and WMMA atom in each target's catalog, plus the 4x4x4 atoms, through both engines and requires the same source or the same refusal. It also checks that on gfx1250 the s_waitcnt / sync_lds_only refusal has the same type and exact message in both engines, and that the other targets still lower. The C++ inline-asm port refused malformed ops that Python failed on with a bare KeyError / IndexError. Both engines now give the same typed error for these, and the test checks each one.
  • The parity batch worker no longer keeps its own list of emitter modes. It passes the mode through, and the emitter accepts it or answers "unknown mode".

Results

  • Rebased on develop.
  • ctest 36/36.
  • Byte identity is green at llvm20, llvm22, llvm23 and llvm24. No golden file changed.
  • The HIP source identity test passes for every family.
  • Full platform pytest, both default and both-engines strict, fails only tests that also fail on develop.

@cgmillette

Copy link
Copy Markdown
Contributor Author

Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max

PR #14845 — make both engines emit the same HIP source (cgmillette, AICK-2903) — 🟠 Follow-up: N1 fixed, ABI bump owed

Prior item Status
Our 11-02 N1: C++ float spelling not repr at powers of two 🟢 FIXED e7def4d8
• probe: 2^-24 and all 2098 powers of two match (ran)
• 123k f16/f32/f64 constants and fills: same text (ran)
shumway (11-02) SRP/DRY ×4: formatter home, copies, focused test, locale 🟢 FIXED e7def4d8
• one primitive in float_format.h; both copies gone (read)
• table ctest + 150k sweep pass; de_DE locale too (ran)
• table ctest: superbuilds only; tier YAML lacks it (read)
shumway (11-02) APPROVED 🟢 APPROVED 8cc38cea
• older head; still counts on develop
• rocke-lib owner still needed: library/tests/parity/ (read)
CongMa13 (11-02) reply on shumway's comment, resolve threads 🟢 Answered
• cgmillette replied 11-02 20:27Z (read)
• no review threads exist: 0 of 0
Merge with develop (14868 rewrote the test) 🟡 OPEN
• conflict: test_engine_linkage.py only (ran)
• keep develop's side; add "ROCKE_FLOAT_FORMAT_TEST": native_artifact("rocke_float_format"),
• resolved: contract test 2/2; head's failure is develop's 14875 (ran)
Bots / CI 🟡 Not run
• head conflicts: no build or test workflow ran (read)
• 3 ✓ · 1 cancelled bot run; Math CI status ✓
• merge: dirty, 1 file
Our checks 🟢 Clean on the resolved merge
• merge vs cee1e23b + the resolution above (ran)
• build + ctest 35 ✓ / 1 skip; byte identity every family ×4 GREEN; goldens, differential unchanged (ran)
• PR tests 310 ✓ / 1 skip in 3 modes; HIP route 3113 ✓; rules 1–3 clean (ran)
Code quality 🟠 Should-fix
• formatter, literal adapters, MMA tables: one home each (read)
• tests: N3 copies engines.py; N4 WMMA operands Python-only
• minor: N5; mode list spelled ×9; C++ HIP checks have no CI lane
# Sev file:line Finding
N2 🟠 include/rocke/lower_hip.h:54-63, include/rocke/abi.h:93, portable_ir/src/abi.py:78 rocke_hip_arch_t drops two fields; no ABI bump
• impact: 24 → 16 bytes, returned by value; a C caller built on develop's header reads garbage (ran)
• fix: bump ROCKE_ABI_VERSION and BINARY_ABI with a history line; develop is 6; takes develop + 1 when it lands
N3 🟠 tests/core/test_hip_lowering_parity.py:21-47, :127; differential/test_hip_source_identity.py:57, :121, :153 New tests rebuild engines.py / run_diff.py helpers
• impact: two homes for one both-engine check; env copy skips the ROCKE_DEBUG_LOC pop; stderr-wide "unknown mode" match
• fix: lower_hip_with_both_engines / assert_both_engines_reject_hip (args[0] and status there), _python_emitter_env, mode_unsupported, known_arches()
N4 🟠 tests/core/test_mma_operand_register_type.py:15, :298, :419-459 Ported WMMA operand rules tested through Python only
• impact: C++ matches today (ran: 3 tests, 5 refusals); no test holds it
• fix: ported op_ids into _NATIVE_HIP_MMAS; both-engine helper in the 3 HIP tests; correct the docstring
N5 🟡 core/lower_hip.py:1509-1513, lower_hip/lower_hip_mma.cpp:648 A non-string template is reported as missing
• impact: template=7 raises KeyError "missing 'template' attr"; develop raised TypeError (ran)
• fix: absent → KeyError; not a string → TypeError naming the type, both engines

Utility: 🟢 Positive — both engines emit identical HIP for every parity kernel; repr-exact floats.


Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted

Follow-up review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: range-diff, merge with develop (conflict resolved), native build + ctest, byte identity and goldens at four LLVM flavors, differential, the PR's tests in three modes, HIP route tests and the identity gate with the engine, both-engine float and MMA sweeps, ABI layout, CI at head, code quality. Flags whether a closer look is needed; not an approval or rejection.

@CongMa13 CongMa13 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

One new issue found; see the inline comment about multiline inline-assembly templates producing an unescaped newline in the generated C++ string literal.

@cgmillette
cgmillette force-pushed the users/cgmillette/rocke/hip-text-identity branch from 83770a6 to dd6473f Compare November 3, 2026 22:02
@cgmillette
cgmillette requested a review from a team as a code owner November 3, 2026 22:02
@cgmillette

Copy link
Copy Markdown
Contributor Author

Rebased on develop and addressed the open review feedback. This is a summary; the replies on each item have the details.

  • Rebase: one conflict in tests/test_engine_linkage.py. I kept fix(rocke): scope installed compiler test policy #14868's installed-test policy and added ROCKE_FLOAT_FORMAT_TEST through its native_artifact helper. rocke_float_format is now in the quick and ffm-quick tiers of ROCKE_ENGINE_test_categories_external.yaml.
  • Unsigned integer WMMA rows: develop added these rows, and the C++ HIP lowering passed a fixed true, ..., true for their signedness flags. Both lowerings, and the LLVM matrix-call lowering, now read the flags from the catalog row through rocke_matrix_operand_is_signed. The per-atom parity test builds every new operand format, including fp6, fp4, u8 and u4.
  • Inline asm escaping (review comment): the template, the constraints and the clobber are escaped for a C++ string literal in both engines, so a multi-line template compiles.
  • Inline asm refusals: a template or constraints attr that is not a string is now a TypeError in both engines.
  • ABI: ROCKE_ABI_VERSION and BINARY_ABI are bumped to 8 for the smaller rocke_hip_arch_t.
  • Tests: the shared both-engine HIP helpers now compare the error type too. The identity gate requires matching refusals; that change exposed one RuntimeError/ValueError mismatch, which is fixed in Python.

Local gates on the head:

  • Build and ctest: 47/47 pass.
  • Byte identity at the host LLVM flavor: GREEN, 90 + 6 families. LLVM output is not meant to change.
  • The PR's tests pass with ROCKE_BACKEND=both strict.
  • Platform suite with ROCKE_BACKEND=both strict: one failure, test_the_cpp_fallback_does_not_warn_twice, which also fails on develop.
  • pre-commit is clean.

@cgmillette

Copy link
Copy Markdown
Contributor Author

Re our review (#14845 (comment)):

Follow-up on N2-N5 (rebased on develop, head dd6473f):

  • N2 (ABI): ROCKE_ABI_VERSION and BINARY_ABI go to 8, with a history line for the smaller rocke_hip_arch_t. The bump is in the commit that removes the fields (4644b6b).
  • N3 (test helpers): the HIP tests now use rocke.tests.engines. hip_outcomes_of_both_engines returns each engine's source or its refusal as a HipRefusal (error type and message), and both reject helpers compare the Python message by its first argument, so a KeyError compares as written. test_hip_source_identity.py takes run_diff.python_emitter_env, run_diff.mode_unsupported and a new run_diff.emission_from_record, and now requires the same refusal (type and message), not just two refusals. That turned up one real difference: the Python HIP lowering raised RuntimeError for an LDS access before its allocation, where C++ reports ValueError. Python now raises ValueError (2fdd727). The MMA atom cases cover every target with its own catalog, which adds gfx908.
  • N4 (WMMA operands): the integer and fp8 WMMA HIP cases in test_mma_operand_register_type.py and test_integer_wmma_signedness.py run both engines. The docstring no longer says C++ lacks those lowerings.
  • N5 (template type): a template or constraints attr that is not a string is now a TypeError in both engines. A missing one is still a KeyError.

The rebase also needed one C++ fix: develop added the unsigned integer WMMA rows, and the C++ HIP lowering hard-coded signed flags. Both lowerings now read the flags from the catalog row through one primitive, rocke_matrix_operand_is_signed.

@cgmillette

Copy link
Copy Markdown
Contributor Author

Automated overlap notice (rocKE merge monitor)

This PR and #12167 both add the HIP lowering of tile.inline_asm in both engines. Each one replaces develop's C++ refusal (lower_hip_mma.cpp:354) and rewrites the Python _op_tile_inline_asm:

  • here: h_asm_operands and rocke_h_op_tile_inline_asm (lower_hip_mma.cpp:608, :657), the escaper rocke_h_escape_cxx_string (core.cpp:601), and _escape_cxx_string and _op_tile_inline_asm (lower_hip.py:324, :1542);
  • in feat(rocke): add NVFP4 GEMM with tensor scales and focused example #12167: h_escape_asm_string, h_parse_asm_constraints, h_asm_operands and rocke_h_op_tile_inline_asm (lower_hip_mma.cpp:348, :442, :534, :559), and _parse_hip_asm_constraints and _op_tile_inline_asm (lower_hip.py:103, :1583).

They differ on mismatched operands: this PR pairs constraints with operands "stopping at the shorter list" and truncates (lower_hip_mma.cpp:606, lower_hip.py:1587), while #12167 refuses a count mismatch (lower_hip_mma.cpp:518-529).

Owner: #12167 keeps it (cgmillette's priority, and it lands first). Resolution: this PR drops its inline-asm lowering and its escaper in both engines and keeps the rest of its HIP parity work on top of #12167's lowering.

Suggested order: #12167 lands first, then this PR rebases onto develop, then #14881, which is stacked on this PR, rebases onto it. The merge plan already has #12167 in its first group with auto-merge on, and it has this PR wait for #12167 because both change lower_hip.py and lower_hip_mma.cpp.

Rebase: after #12167 lands:

Then run that file with tests/core/test_optimization_barrier.py and the rocke_optimization_barrier ctest.

If you think this PR should own it instead, reply here and on #12167, and the reviews will follow the new decision.

Automated overlap check: one owner per shared piece of code, coordinated on the PRs. Nothing else changed.

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

The integer and FP8 WMMA emitters access results and operands without arity checks, allowing malformed serialized IR to crash the host process.

5 open issues
What changed in this PR

Aligns rocKE’s Python and C++ HIP-source lowerers and extends the existing engine-parity checks to cover HIP output.

Changes:

  • Centralizes Python-compatible float formatting and escapes inline assembly strings.
  • Adds missing C++ MMA and inline assembly lowerings, aligns target refusals, and updates the binary ABI.
  • Adds HIP-source identity tests, target-tagged emitter output, and float-format test wiring.

Paths are relative to dnn-providers/hip-kernel-provider/.

File Description
rocke/​TESTING.md Documents the HIP-source identity test.
rocke/​platform/​tests/​tools/​test_run_all.py Checks float-format fixture discovery and propagation.
rocke/​platform/​tests/​test_engine_linkage.py Adds the installed float-format fixture expectation.
rocke/​platform/​tests/​run_all.py Registers the native float-format fixture.
rocke/​platform/​tests/​README.md Documents float-format parity execution.
rocke/​platform/​tests/​instances/​parity/​moe_sorting_emit.py Adds target-tagged IR output.
rocke/​platform/​tests/​instances/​parity/​moe_fused_mega_fp8_emit.py Adds the shared target-ir mode.
rocke/​platform/​tests/​instances/​parity/​layernorm2d_emit.py Adds the shared target-ir mode.
rocke/​platform/​tests/​instances/​parity/​ir_serialize_emit.py Supports target-tagged serialization.
rocke/​platform/​tests/​instances/​parity/​grouped_gemm_stress_emit.py Adds target-tagged grouped GEMM output.
rocke/​platform/​tests/​instances/​parity/​gemm_stress_emit.py Adds target-tagged GEMM output.
rocke/​platform/​tests/​instances/​parity/​fused_moe_emit.py Recognizes and emits target-tagged IR.
rocke/​platform/​tests/​instances/​parity/​fused_moe_e2e_emit.py Emits one target-tagged block per stage.
rocke/​platform/​tests/​instances/​parity/​batched_gemm_stress_emit.py Adds target-tagged batched GEMM output.
rocke/​platform/​tests/​instances/​parity/​add_rmsnorm2d_rdquant_emit.py Adds the shared target-ir mode.
rocke/​platform/​tests/​instances/​parity/​_emit_common.py Centralizes emitter modes and target tagging.
rocke/​platform/​tests/​instances/​differential/​test_hip_source_identity.py Compares corpus-wide HIP source and refusals.
rocke/​platform/​tests/​instances/​differential/​test_gfx950_scaled_mfma_emit.py Uses the public emitter-environment helper.
rocke/​platform/​tests/​instances/​differential/​run_diff.py Exposes environment and emission-record helpers.
rocke/​platform/​tests/​instances/​differential/​python_batch_worker.py Delegates mode validation to emitters.
rocke/​platform/​tests/​core/​test_mma_operand_register_type.py Extends HIP operand checks to both engines.
rocke/​platform/​tests/​core/​test_integer_wmma_signedness.py Checks both HIP engines’ signedness flags.
rocke/​platform/​tests/​core/​test_hip_lowering_parity.py Adds atom, inline assembly, and refusal parity cases.
rocke/​platform/​tests/​core/​test_float_format.py Compares native formatting with Python repr.
rocke/​platform/​tests/​core/​test_float_format.cpp Adds formatter cases, locale checks, and sampling support.
rocke/​platform/​tests/​core/​matrix_call_lowering.cpp Tests the shared signedness primitive.
rocke/​platform/​tests/​cmake/​CoreTests.cmake Builds, registers, and installs the formatter test.
rocke/​platform/​python/​rocke/​tests/​engines.py Adds comparable HIP-source and refusal outcomes.
rocke/​platform/​python/​rocke/​portable_ir/​src/​abi.py Updates the Python binary ABI to 8.
rocke/​platform/​python/​rocke/​core/​lower_hip.py Aligns literals, escaping, validation, and refusals.
rocke/​platform/​cpp/​support/​float_format.cpp Implements Python-compatible float formatting.
rocke/​platform/​cpp/​include/​rocke/​matrix_calls_internal.h Declares shared signedness and integer-row helpers.
rocke/​platform/​cpp/​include/​rocke/​lower_hip.h Updates the HIP architecture representation.
rocke/​platform/​cpp/​include/​rocke/​lower_hip_internal.h Updates literal, escaping, and target-gate interfaces.
rocke/​platform/​cpp/​include/​rocke/​ir_serialize.h Documents shared float formatting.
rocke/​platform/​cpp/​include/​rocke/​float_format.h Declares the common float formatter.
rocke/​platform/​cpp/​include/​rocke/​abi.h Records the layout change and ABI version 8.
rocke/​platform/​cpp/​core/​lower_llvm/​matrix_call.cpp Reuses the shared signedness primitive.
rocke/​platform/​cpp/​core/​lower_hip/​lower_hip_mma.cpp Adds MMA and inline assembly emitters.
rocke/​platform/​cpp/​core/​lower_hip/​lower_hip_mem.cpp Updates the LDS-refusal comment only.
rocke/​platform/​cpp/​core/​lower_hip/​core.cpp Adds shared formatting, escaping, and target checks.
rocke/​platform/​cpp/​core/​isa/​matrix_calls.cpp Centralizes signedness derivation and integer rows.
rocke/​platform/​cpp/​core/​ir/​serialize.cpp Replaces the duplicated float formatter.
rocke/​platform/​cpp/​core/​ir/​print.cpp Uses the common formatter for float attributes.
rocke/​platform/​CMakeLists.txt Propagates the installed formatter-test path.
rocke/​library/​tests/​parity/​gfx950_attention_tiled_3d_emit.py Tags segment and reduction kernels with their target.
rocke/​library/​tests/​parity/​gfx942_attention_tiled_3d_emit.py Adds target-tagged attention output.
rocke/​library/​tests/​parity/​fmha_splitkv_decode_emit.py Adds target-tagged decode output.
rocke/​library/​tests/​parity/​conv_direct_grouped_emit.py Emits each convolution kernel’s target.
rocke/​library/​tests/​parity/​_emit_common.py Adds shared target-tagged library emission.
ROCKE_ENGINE_test_categories_external.yaml Enrolls the formatter test in quick tiers.

Review effort: Max

cgmillette and others added 3 commits November 4, 2026 17:53
The C++ HIP lowering printed float constants with %g, which keeps only six
significant digits and drops the decimal point from whole numbers. That
emitted invalid C++ ("0f", "1024f") and changed values ("-1.4427f" for
-1.4426950408889634), so the C++ HIP source differed from the Python
engine's and did not compile.

The Python engine writes every float with repr(): in printed IR, in
serialized IR and in HIP literals. The C++ engine now writes all three with
one formatter, rocke_append_python_float_repr in support/float_format.cpp,
which replaces the two copies the IR printer and the serializer each kept.
It takes the shortest round-trip digits from std::to_chars and lays them out
the way CPython does. The copies it replaces probed printf precisions until
strtod read the value back; that picked 17 digits for some values that have
a 16-digit spelling (2**-24 printed as 5.9604644775390625e-08, repr is
5.960464477539063e-08) and depended on the C locale's decimal point.

Each engine's HIP lowering adds one thin adapter on top: repr() plus an "f"
suffix for a float literal, repr() alone for a double literal, and the
<math.h> NAN / INFINITY macros cast to the literal's type for the values repr()
spells in a way C++ does not accept. The float and double entry points are
one-line calls into that adapter.

rocke_float_format checks the formatter against a table of values whose repr
is known (zeros, integral values, both sides of the 1e-4 and 1e16 notation
switches, 2**-24 and its neighbours, subnormals, the finite extremes, nan
and the infinities), again under a comma-decimal locale when one is
installed. test_float_format.py compares it with repr() over every power of
two and of ten, the doubles either side of each, and a fixed random sample,
and lowers float and double constants (zeros, integral values, the notation
switches, 2**-24, the finite extremes, nan and the infinities) through both
engines' HIP lowerers and checks every literal. tests/run_all.py and the
installed rocke_pytest point it at the rocke_float_format binary
(ROCKE_FLOAT_FORMAT_TEST), as they do for the other native parity tests.
The ctest is in the quick and ffm-quick tiers of
ROCKE_ENGINE_test_categories_external.yaml, beside rocke_ir_types.

Co-Authored-By: Claude <noreply@anthropic.com>
A newer target replaced the monolithic s_waitcnt with split wait counters,
and clang cannot select __builtin_amdgcn_s_waitcnt there. The two HIP
lowerings still emitted it, each with a different immediate: Python used
the gfx9 layout, C++ the gfx11 layout.

Both now ask the target's ISA backend whether it emits the monolithic
s_waitcnt (the fact the LLVM lowering already uses) and refuse with the
same NotImplementedError message when it does not. The C++ HIP arch seam
records this as a new waitcnt family, ROCKE_HIP_WAITCNT_NONE. Every other
target keeps its current encoding.

test_hip_lowering_parity.py lowers tile.s_waitcnt and tile.sync_lds_only
through both engines: on the target with split wait counters both refuse
with that exact message (C++ with ROCKE_ERR_NOTIMPL), and on gfx942, gfx950, gfx1151 and gfx1201 both
emit the same source.

Co-Authored-By: Claude <noreply@anthropic.com>
The C++ HIP lowering refused kernels the Python one lowers: every f16,
bf16, fp8 and bf8 MFMA atom outside the f32/f64 controls, and the gfx11,
gfx12 and newer WMMA atoms all failed with "no HIP lowering for op".

Each is now ported from its Python handler and emits the same text. The
atoms are grouped by how their builtin takes its operands, with one table of
op_id-to-builtin names and one small emitter per group:

- MFMA f16/bf16 atoms, including the 4x4x4 ones, pass A, B and C to their
  builtin as they are;
- fp8/bf8 MFMA atoms copy the packed operands into long scalars first;
- the gfx11 and gfx12 f16 WMMA atoms pass A, B and C as they are, and a newer
  target's K=32 ones use the builtin's 8-operand form;
- a newer target's fp8/bf8 WMMA atoms and the iu8/iu4 ones take A and B as packed i32
  vectors. The iu8/iu4 atoms lower under every signedness row of the
  catalog, and their builtin's signedness bools come from the row through
  rocke_matrix_operand_is_signed, which the LLVM lowering's flags now read
  too (Python's Signedness.flag is the twin). ROCKE_INTEGER_MMA_ROWS lists a
  table's rows for each signedness, as integer_mma_op_ids does in Python.

The accumulator width of each atom, and the packed word count of A and B, come
from the arch MMA catalog (c_frag_len, a_frag_len, b_frag_len) rather than a
second table.

rocke_h_require_wmma_arch now admits an atom by the target's MMA catalog,
as Python's _require_wmma_arch does. The has_wmma and family fields of
rocke_hip_arch_t, which it used to read, are removed. That changes the
struct's size, so ROCKE_ABI_VERSION and the Python BINARY_ABI go to 9.

Both HIP lowerings of tile.inline_asm now read its template and
constraints attrs through one helper each (_inline_asm_string_attr,
h_inline_asm_string_attr), so a serialized kernel that lacks one is
refused with the same KeyError, and a non-string one with the same
TypeError, by both engines. Python used to raise a bare KeyError for a
missing template and a TypeError from re.sub for a non-string one.

test_hip_lowering_parity.py lowers every native MFMA and WMMA atom in each
target's catalog and the 4x4x4 atoms through both engines and requires
the same source or the same refusal, and checks the type and message of
each template and constraints refusal. It builds each A and B operand from zero words,
so the formats with no zero constant (the 6- and 4-bit floats, the unsigned
integers) are lowered too. The HIP cases of test_integer_wmma_signedness.py
and test_mma_operand_register_type.py, which ran the Python lowering alone
for the atoms the C++ one did not have, now run both engines, and the check
that no engine writes a signedness flag as a literal also reads the C++ HIP
format strings.

Co-Authored-By: Claude <noreply@anthropic.com>
cgmillette and others added 6 commits November 4, 2026 17:57
The LLVM-IR route is held byte-identical between the engines; the HIP
source route was not, and drifted. This adds a test that lowers every
parity kernel through both HIP lowerers (Python lower_kernel_to_hip and
rocke_engine.lower_serialized_ir_hip) and requires the same text, or a
refusal from both. It is parametrized by parity family, covers the
platform and library parity trees, and skips without rocke_engine.

The serialized IR does not record the target a kernel is built for, so
the parity emitters gain a "target-ir" mode: a "target <gfx>" line, then
the serialization, once per kernel. run_emit accepts it, and every emitter
that parses its own argv does too, except conv_abi_emit.py, which builds no
kernel. The batch worker no longer keeps its own list of modes: it passes
the mode to the emitter, which accepts it or answers "unknown mode". The
test fails if an emitter does not accept the mode or builds no kernel.

Co-Authored-By: Claude <noreply@anthropic.com>
…r in HIP

The Python HIP lowering raised RuntimeError when an op reads or writes
an LDS allocation whose smem_alloc has not been lowered yet; the C++ HIP
lowering refuses the same kernel with the same message and status 1,
which stands for ValueError (there is no RuntimeError status). The
Python lowering now raises ValueError at each of those eleven sites, so
both engines refuse such a kernel with the same error type.

Co-Authored-By: Claude <noreply@anthropic.com>
The HIP lowering tests compared the two engines through helpers of their
own. They now use the ones in rocke.tests.engines:

- assert_both_engines_reject and assert_both_engines_reject_hip compare
  the Python message by its first argument, so a KeyError refusal
  compares as written rather than quoted.
- hip_outcomes_of_both_engines returns each engine's HIP source for a
  serialized kernel, or its refusal as a HipRefusal: the error type (the
  Python exception's, or the one the C++ status code stands for) and the
  message. test_hip_lowering_parity.py requires equal outcomes for every
  MMA atom and checks each refusal it pins (s_waitcnt, malformed inline
  asm) against a HipRefusal, so the engines must agree on the error type
  as well as the message.
- test_hip_source_identity.py requires equal outcomes too: the same
  source, or the same refusal, where it used to accept any two refusals.

test_hip_source_identity.py also takes the emitter environment from
run_diff.python_emitter_env (public now, so the debug-location capture
stays off as it does for the gate), detects an emitter refusing the
target-ir mode with run_diff.mode_unsupported, which reads only the last
line of its stderr, and reads each worker record with the new
run_diff.emission_from_record, which the gate's batch reader now uses
as well.

The MMA atom cases now cover each target with its own catalog
(arches_with_own_facts): gfx908 is added, and the gfx11 and gfx12 atoms
run on gfx11-generic and gfx12-generic, whose catalogs gfx1151 and
gfx1201 take, instead of on those two.

Co-Authored-By: Claude <noreply@anthropic.com>
A kernel read from serialized IR is not verified, so a tile.mma op can
arrive without its accumulator or its result. The C++ HIP lowering of the
fp8 and integer WMMA atoms, the f32 and f64 MFMA controls and the TF32
atoms indexed both regardless and could read past the op's arrays; the
Python lowering failed with whatever unpacking or indexing error came
first. A scaled WMMA op without its two scales was refused by C++ with
"<op_id> expects 5 operands and 1 result" and by Python with an unpacking
error.

Every path of both lowerings that lowers an atom from A, B and C now
checks first that the op has those three operands and one result, and
refuses it otherwise with the same ValueError: "tile.mma <op_id> expects
3 operands and 1 result" (h_require_abc, Python's _require_abc_atom). An
op_id neither engine lowers is still refused as unknown, before any count
is checked. The Python scaled WMMA lowering checks for five operands with
the C++ message.

test_hip_lowering_parity.py drops the accumulator, and separately the
result, from a kernel for every native atom in each target's catalog and
requires both engines to refuse it alike, with that message wherever the
well-formed kernel lowers. It also gives every scaled atom an op with A,
B and C only, and requires the same refusal from both, the five-operand
one for the scaled WMMA atoms.

Co-Authored-By: Claude <noreply@anthropic.com>
The test compared the kernels each config printed and ignored how the
config's emitter run ended. A config that failed after printing part of
its output was compared as if it had built, and one that failed with no
output dropped out of the comparison, so a broken emitter could remove
kernels from the gate unnoticed.

Each config now passes only if it printed its kernels and exited 0, or
reached the end of the config list. A config that fails with no output
passes only as the builder refusing it: it must have raised one of the
exception types that stand for a refusal (REFUSAL_ERROR_TYPES in
rocke.tests.engines: ValueError, TypeError, KeyError and
NotImplementedError, the types the C++ engine's status codes stand for),
and the ir mode, which builds the same kernel and only serializes it,
must refuse it with the same error. A config that exits non-zero after
printing, exits 0 without printing, fails with another exception type,
fails only in the target-ir mode, or fails in the ir mode with another
error fails the family.

run_diff.py makes public the two facts the test reads from a config's
run: is_end_emission (was the index past the last config) and
last_error_line (the last stderr line, which mode_unsupported also
reads). test_hip_source_identity.py checks the classification against an
emitter with one config for each way a run can end; that case needs no
engine.

Co-Authored-By: Claude <noreply@anthropic.com>
The test lowers each kernel through the C++ engine's rocke_engine
binding and skipped without it. CI runs CTest against the installed
artifact, whose wheel has no binding, and the test was not installed, so
it never ran there.

ir_lower_cli.cpp, the tool that lowers serialized IR with the C++
engine, gains a --hip option: it lowers with rocke_lower_kernel_to_hip
and the defaults the binding uses, writes the source unchanged (all three
streams binary on Windows), and on a refusal writes the engine's status
and message in the form the binding's errors take. CMake now builds it as
rocke_ir_lower_cli and installs it beside the test binaries, outside the
list that registers each installed binary as a test.

In rocke.tests.engines, hip_outcomes_of_both_engines runs the C++ HIP
lowering through the binding when it is built and otherwise through the
program ROCKE_IR_LOWER_CLI names (ir_lower_cli_hip_outcome);
skip_without_native_hip_lowering skips a test that has neither, and a
name that is not a file fails. test_hip_lowering_parity.py requires the
program and the binding to report the same source and the same refusal
where both are present.

The test is installed with the differential files the golden checks
already install, and the rocke_hip_source_identity CTest runs it with
ROCKE_IR_LOWER_CLI set, a 120 s limit per family and 900 s in all; one
family's emitter run is bounded at 100 s, below the per-family limit. It
is listed in the quick and ffm-quick tiers. The rocke_pytest CTest does
not collect it.

Co-Authored-By: Claude <noreply@anthropic.com>
@cgmillette
cgmillette force-pushed the users/cgmillette/rocke/hip-text-identity branch from dd6473f to 92241b6 Compare November 4, 2026 23:59
@cgmillette

Copy link
Copy Markdown
Contributor Author

Rebased onto current develop and addressed the open review items. Summary of what changed since the last push:

Rebase. #12167 landed with its own HIP lowering of tile.inline_asm in both engines, so this PR no longer ports or escapes inline asm: it takes develop's lowering and drops its own escaper (_escape_cxx_string, rocke_h_escape_cxx_string) and constraint pairing. What it keeps for inline asm is one helper per engine for the template and constraints reads, so a serialized kernel missing one is refused with the same KeyError, and a non-string one with the same TypeError, by both engines. The rocke_hip_arch_t size change now takes ROCKE_ABI_VERSION / BINARY_ABI 9 (develop is 8).

MMA arity (Copilot, the fp8 and integer WMMA emitters). Every path of both HIP lowerings that lowers an MMA atom from A, B and C now checks first that the op has three operands and one result, and refuses it with the same ValueError otherwise. That covers the fp8 and integer WMMA emitters and also the f32/f64 MFMA controls and TF32 atoms, which indexed the same arrays unchecked. The Python scaled WMMA lowering now counts its five operands with the C++ message. test_hip_lowering_parity.py drops the accumulator, and separately the result, for every native atom in every catalog, and gives every scaled atom an op without its scales; both engines must refuse alike.

sideeffect (Copilot). Already fixed on develop by #12167: both engines refuse a present non-boolean sideeffect (including "") with "HIP inline asm sideeffect must be a boolean" and default to true only when it is absent. test_optimization_barrier.py::test_hip_native_rejects_malformed_barrier_attributes checks both engines with "", "false", 0, 1, 0.0 and 1.0.

Identity gate in CI (Copilot). CI runs installed CTest without the rocke_engine binding. tests/core/ir_lower_cli.cpp gains --hip and is built and installed as rocke_ir_lower_cli; rocke.tests.engines lowers through it when the binding is absent. The new rocke_hip_source_identity CTest runs the gate in the installed tree (about 2.5 minutes locally, 120 s per family, 900 s in all) and is listed in the quick and ffm-quick tiers. A test holds the program and the binding to the same output where both exist.

Emission status (Copilot). The gate now fails a config that exits non-zero after printing, exits 0 without printing, or fails with an exception other than a refusal type (ValueError, TypeError, KeyError, NotImplementedError). A config that fails with no output counts as the builder refusing it only if the ir mode refuses it with the same error. A small test covers each way a config can end.

Our follow-up review (11-03). The rocke_hip_arch_t ABI bump, the shared both-engine helpers and the both-engine WMMA operand tests were already in the previous push (the commits that lower the MMA atoms and that share the helpers). The non-string template is a TypeError again, in both engines.

CI. The red Test hipkernelprovider shards (develop's fp64 doc test, fixed by #15002) and Stage - Math Libs (develop's hipCCL change, reverted by #15089) are not from this PR; this push picks up both fixes. The other failures were runner infrastructure.

Checked locally: build and ctest, byte identity at llvm20/22/23/24 GREEN, the identity gate with the binding and from an installed tree through rocke_ir_lower_cli (100 passed), the PR's tests with both engines in strict mode, the platform suite (one failure, test_the_cpp_fallback_does_not_warn_twice, which also failed on develop in that mode when we last compared), pre-commit over the PR range.

cgmillette and others added 3 commits November 4, 2026 22:30
The C++ parser of serialized IR kept the names of the values it had
defined in an array and found each operand by comparing it with every
name defined before it, so parsing a kernel took time quadratic in its
size. For a large attention kernel, almost all of the time the C++ HIP
lowering took, through rocke_engine.lower_serialized_ir_hip or the
serialized-IR lowering entry points, went to that search.

The table is now an open-addressing hash table on the name (64-bit
FNV-1a, linear probing, kept at most three-quarters full), in the
builder's arena as before. It keeps the same contract: a later definition
of a name replaces the earlier one, and the names a loop defines are
removed when the loop ends, so a sibling loop can define them again.

test_serialized_ir_name_scopes.py builds a kernel of 200 sibling loops
that each define the same induction variable and iter-arg names, with
over a thousand values in all, and requires both engines to lower it to
the same LLVM IR and the same HIP source.

Co-Authored-By: Claude <noreply@anthropic.com>
…test

The HIP source identity test ran the C++ lowering of every kernel on its
own: one rocke_ir_lower_cli process per kernel where the binding is
absent, and it ran each family's emitter only when that family's test
started.

- rocke_ir_lower_cli's --hip option becomes --hip-batch, which lowers a
  list of kernels in one run: frames of "<arch> <byte count>" and the
  serialized IR in, frames of "<status> <byte count>" and the source or
  the engine's refusal out, with the binding's fallback messages. An empty
  or malformed frame, or a reply that cannot be written, ends the run with
  status 2, as the binding refuses empty IR outright. All three streams
  are binary on Windows.
- In rocke.tests.engines, native_hip_outcomes lowers a list of kernels
  through the binding or one program run (ir_lower_cli_hip_outcomes, which
  checks every reply frame), python_hip_outcome is the Python half, and
  hip_outcomes_of_both_engines is built from the two.
- The test starts each selected family's emitter run ahead of the test that
  reads it, in the session's order and a few at a time, so the emitters
  overlap the Python lowering of the families before them. It starts none
  where the C++ lowering cannot run, and cancels the runs not yet started
  when the module ends.
- The tie test between the program and the binding sends one batch that
  holds source and refusals.

The rocke_hip_source_identity CTest limit drops from 900 s to 300 s.

Co-Authored-By: Claude <noreply@anthropic.com>
The Windows hipkernelprovider CI shards run close to their 30-minute
step limit, and rocke_hip_source_identity pushed one of them over it.
The CTest is no longer registered on Windows, as rocke_library_gdn_pytest
is not; it stays registered, in the quick and ffm-quick categories, on
every other platform, so the Linux hipkernelprovider jobs still run it on
every pull request.

The cost is that CI no longer compares the MSVC-built engine's HIP
output with Python's. The CMake comment, the category list and
TESTING.md say so, and say to register it again once the Windows shards
have room.

Co-Authored-By: Claude <noreply@anthropic.com>
@cgmillette

Copy link
Copy Markdown
Contributor Author

The new rocke_hip_source_identity CTest made the Windows hipkernelprovider shards too slow. It ran for about 219 s in each shard, and shard 2 hit the 30-minute step limit. Three new commits fix this.

  • perf(rocke): look up SSA names by hash when parsing serialized IR. Most of the gate's C++ time was not spent in the HIP lowering. It was spent in the C++ parser of serialized IR, which looked up every operand name by comparing it with every name defined before it. That is quadratic on the large attention kernels. The parser now uses a hash table, so every C++ path that lowers serialized IR is faster, the binding included. Parsed kernels are unchanged: byte identity is GREEN at all four LLVM flavors. A new test holds both engines to the same output for a kernel whose 200 sibling loops reuse their names.
  • test(rocke): lower a family's kernels in one run in the HIP identity test. rocke_ir_lower_cli --hip-batch now lowers all of a family's kernels in one process. The test also starts each family's emitter run ahead of the test that reads it. The whole gate now takes about 47 s locally, down from about 150 s, and its CTest limit is 300 s.
  • test(rocke): do not register the HIP source identity CTest on Windows. This follows the precedent of rocke_library_gdn_pytest. The cost is that CI no longer compares the MSVC-built engine's HIP output with Python's. It matched on every parity kernel on both Windows lanes in the last run. Float formatting, the most host-sensitive part, keeps its own rocke_float_format test on Windows. The CMake comment says to register it there again once those shards have room.

Every lane runs ctest -L '^standard$', and the quick category carries the standard label. So which lanes run the test depends on where it is registered, not on its tier:

Lane Before After
Windows gfx110X, both shards about 219 s each not registered: 0 s
Linux hipkernelprovider shards the full gate, about 150 s locally the full gate, about 47 s locally (expect roughly 50 to 75 s on CI)

The test stays in the quick and ffm-quick categories, so every Linux hipkernelprovider job runs it on every PR.

Checked locally:

  • build and ctest;
  • byte identity GREEN at llvm20/22/23/24;
  • from an installed tree, ctest -R rocke_hip_source_identity: 100 passed in 47 s;
  • the PR's tests with both engines in strict mode;
  • platform suite: one failure, test_the_cpp_fallback_does_not_warn_twice, which also failed on develop in this mode when we last compared;
  • pre-commit over the PR range.

@cgmillette

Copy link
Copy Markdown
Contributor Author

Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max

PR #14845 — make both engines emit the same HIP source (cgmillette, AICK-2903) — 🟠 Follow-up: rebase owed, CI coverage gap

Prior item Status
Our 11-03 N2: rocke_hip_arch_t drops two fields, no ABI bump 🟢 FIXED f57fae89
• ROCKE_ABI_VERSION and BINARY_ABI 9, with a history line; probe FIXED (ran)
• develop is 8; takes develop + 1 when it lands
Our 11-03 N3: new tests rebuild engines.py / run_diff.py helpers 🟢 FIXED f97de1c7
• tests take rocke.tests.engines and run_diff helpers (read)
• env copy, stderr-wide match, local status table gone (read)
Our 11-03 N4: ported WMMA operand rules tested through Python only 🟢 FIXED f57fae89
• a C++ and a Python mutant each fail these tests (ran)
• but CI now skips them: N6
Our 11-03 N5: a non-string template reported as missing 🟢 FIXED f57fae89
• missing → KeyError, non-string → TypeError, both engines alike (ran)
Our 11-03 overlap notice: inline asm lowering, with 12167 🟢 RESOLVED
• 12167 landed; the PR dropped its lowering and escapers (read)
• keeps one template/constraints reader per engine (read)
shumway (11-02) APPROVED 🟢 APPROVED 8cc38cea
• older head; still counts on develop
• predates the parser hash table, the CLI and the ABI bump
• code owners still owed: rocKE-lib, hipdnn reviewers (read)
CongMa13 (11-03) multi-line asm template unescaped 🟢 FIXED on develop (12167)
• PR dropped its escaper; develop escapes in both engines (read)
• develop's multi-line template tests pass on the merge (ran)
Copilot (11-04) ×5: WMMA arity ×2, sideeffect, CI lane, emission status 🟢 FIXED 366740e9, db4176ca, 92241b66
• arity checked before every A/B/C read, both engines (ran)
• sideeffect: develop's 12167 refuses non-bools in both (ran)
• installed gate 100 ✓ through the CLI; outcome self-test ✓ (ran)
Merge with develop (14868 rewrote the test) 🟡 OPEN: new conflict
• fu1's test_engine_linkage.py conflict: resolved as advised (read)
• now tests/engines.py ×6 hunks vs 14842 (f28ed8f0): keep develop's helpers (ran)
• re-add the PR's imports, __all__ names, constants and HipRefusal…end, not its _NATIVE_STATUS copy; args[0] into _assert_python_rejects
Bots / CI 🟢 Green (attempt 2)
• attempt 1: Windows hipDNN TestFillTensorWithRandomValues ×2, not rocKE; re-run ✓ (read)
• 45 ✓ · 51 skipped · 1 cancelled; identity CTest 100 ✓ on both Linux shards (read)
• Windows: identity CTest off for shard time, as stated (read); no issue tracks re-adding it
• merge: dirty since 14842 landed; CI saw an older develop
Our checks 🟡 Conflict; resolved here
• merge vs 8b6ada32 + the resolution above: 54 files (ran)
• build + ctest 47 ✓ / 1 skip; byte identity GREEN ×4; goldens ✓; differential unchanged (ran)
• PR tests 2011 ✓ in 3 modes; 36 engines.py users 5710 ✓; installed identity 100 ✓ via the CLI (ran)
• compile 261 configs, 0 new failures; scope and arity probes agree across engines (ran)
Code quality 🟠 Should-fix
• N3, N4 fixed; one home per fact holds (read)
• N6: both-engine HIP tests run in no CI lane (ran)
• minor: N8; TF32 arity calls unreachable (ran); Python reads constraints twice
# Sev file:line Finding
N6 🟠 platform/CMakeLists.txt:662-682, tests/core/test_mma_operand_register_type.py:410, :440 The new both-engine HIP tests run in no CI lane
• impact: CI skips all 698 parity and 37 operand cases; C++ signedness and fp8-operand mutants pass its identity gate (ran)
• our fu1 N4 fix took two operand-cast cases out of CI: develop catches a Python mutant of each, the merge neither (ran)
• fix: ROCKE_IR_LOWER_CLI for rocke_pytest; develop's HIP helpers take the C++ side from native_hip_outcomes
N7 🟠 platform/CMakeLists.txt:628-629, differential/test_hip_source_identity.py:60-61, PR body Two comments and the body state measured times
• impact: platform/AGENTS.md Compliance keeps comments qualitative; a squash commits the body
• fix: relate them to the limits; drop the body's gate and compile timings
N8 🟡 tests/core/ir_lower_cli.cpp:68-112 The CLI repeats the binding's parse-and-lower
• impact: CI's gate trusts the CLI; the test that holds them equal runs in no runner (read)
• fix: one engine entry (parse, lower, message) for both

Utility: 🟢 Positive — repr-exact floats, ported MMA atoms, identity gate in Linux CI.


Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted

Follow-up review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: range-diff, merge with develop (conflict resolved and built), native build + ctest, byte identity and goldens at four LLVM flavors, differential, the PR's tests in three modes, every engines.py user, the identity gate with the engine and installed through the CLI, scope and arity probes in both engines, mutants per engine, hipcc and compile checks, CI at head, code quality. Flags whether a closer look is needed; not an approval or rejection.

@CongMa13 CongMa13 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

One new issue: the HIP source identity gate can pass a family without comparing any successfully emitted source; see the inline comment.

I also agree with two existing open concerns in cgmillette's comment 6655831567: the focused synthetic both-engine HIP cases are skipped in installed CI, and test-duration claims remain in public PR materials.

if difference is not None:
disagreements.append(f"{where}: {difference}")
compared = len(kernels)
assert compared, f"{name}: the emitter built no kernel in any config"

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If both lowerers regress to refusing every kernel in this family with the same error, disagreements stays empty and this check passes: compared counts inputs, not successful source comparisons. Please require at least one successful HIP source match for each expected-lowerable family, with explicit exceptions for intentionally unsupported families, so the parity gate does not pass with no output to compare.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants