Repository navigation
feat(rocke): add FP64 support for elementwise and reduce - #13141
Conversation
github-actions
Bot
added
documentation
project: hip-kernel-provider
rocke: platform
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
e13a9eb to
6e0880c
Compare
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
FP64 reductions mishandle infinite inputs, and HIP lowering has f64 semantic and signed-zero inconsistencies.
Review effort: Balanced
Findings: 1
Open (4)
What changed in this PR
Adds native FP64 support to rocKE’s Python and C++ IR engines, lowerers, memory helpers, and the elementwise/reduce2d kernel families.
Changes:
- Adds f64 types, constants, memory operations, cross-lane operations, and LLVM/HIP lowering.
- Enables native f64 elementwise and reduction kernels.
- Adds parity, build, numeric, dtype, and documentation coverage.
| File | Description |
|---|---|
platform/tests/test_rocke.py |
Adds f64 build and validation tests. |
platform/tests/instances/parity/reduce_emit.py |
Adds Python f64 reduction parity cases. |
platform/tests/instances/parity/reduce_emit.c |
Adds matching C parity cases. |
platform/tests/instances/parity/elementwise_emit.py |
Adds Python f64 elementwise parity cases. |
platform/tests/instances/parity/elementwise_emit.c |
Adds matching C parity cases. |
platform/tests/instances/differential/numeric.py |
Adds f64 GPU differential coverage. |
platform/tests/core/test_dtypes.py |
Tests Python fp64 dtype metadata. |
platform/tests/core/test_dtypes.cpp |
Tests C++ fp64 normalization. |
platform/tests/core/test_amdgpu_lowering_ops.py |
Updates unsupported-type coverage. |
platform/python/rocke/instances/SUPPORT_MATRIX.md |
Documents supported f64 families and ops. |
platform/python/rocke/instances/common/reduce.py |
Implements native f64 reductions. |
platform/python/rocke/instances/common/elementwise.py |
Enables exact f64 elementwise operations. |
platform/python/rocke/helpers/tensor_view.py |
Adds unpromoted vector loading. |
platform/python/rocke/helpers/sweep.py |
Preserves f64 during row sweeps. |
platform/python/rocke/helpers/reduction.py |
Adds f64 LDS reduction support. |
platform/python/rocke/helpers/io.py |
Maps f64 I/O types. |
platform/python/rocke/helpers/distribution.py |
Supports f64 distributed tiles and reductions. |
platform/python/rocke/core/lower_llvm.py |
Lowers f64 types, constants, memory, and min/max. |
platform/python/rocke/core/lower_hip.py |
Adds HIP double lowering. |
platform/python/rocke/core/ir.py |
Defines F64 IR and operations. |
platform/python/rocke/core/dtypes.py |
Adds fp64 aliases and metadata. |
platform/python/rocke/core/__init__.py |
Exports F64 from the core API. |
platform/python/rocke/__init__.py |
Exports F64 publicly. |
platform/dsl_docs/reference/op_vocabulary.md |
Documents f64 operations. |
platform/dsl_docs/primitives/wave_and_cross_lane.md |
Documents f64 shuffling. |
platform/dsl_docs/optimization/optimization_runbook.md |
Adds fp64 validation guidance. |
platform/dsl_docs/ir_lowering/ir_model.md |
Documents the F64 IR model. |
platform/cpp/instances/common/reduce.cpp |
Implements C++ f64 reductions. |
platform/cpp/instances/common/elementwise.cpp |
Implements C++ f64 elementwise kernels. |
platform/cpp/include/rocke/lower_llvm_internal.h |
Declares f64 LLVM formatting. |
platform/cpp/include/rocke/lower_hip_internal.h |
Declares f64 HIP formatting. |
platform/cpp/include/rocke/ir.h |
Exposes C F64 types and constants. |
platform/cpp/include/rocke/instance_reduce.h |
Documents f64 reduction specs. |
platform/cpp/include/rocke/instance_elementwise.h |
Documents f64 elementwise specs. |
platform/cpp/include/rocke/helper_rocke.helpers.tensor_view.h |
Declares scalarized vector loading. |
platform/cpp/include/rocke/helper_rocke.helpers.reduction.h |
Documents f64 reduction input. |
platform/cpp/helpers/tensor_view.cpp |
Implements unpromoted vector loading. |
platform/cpp/helpers/sweep.cpp |
Preserves f64 during C++ sweeps. |
platform/cpp/helpers/reduction.cpp |
Implements f64 LDS reductions. |
platform/cpp/helpers/io.cpp |
Maps f64 I/O types in C++. |
platform/cpp/helpers/distribution.cpp |
Handles f64 distributed LDS reduction. |
platform/cpp/core/lower_llvm/mem.cpp |
Adds 8-byte f64 memory lowering. |
platform/cpp/core/lower_llvm/data.cpp |
Declares f64 min/max intrinsics. |
platform/cpp/core/lower_llvm/core.cpp |
Adds f64 LLVM types, constants, and LDS layout. |
platform/cpp/core/lower_llvm/arith.cpp |
Lowers f64 min/max arithmetic. |
platform/cpp/core/lower_hip/lower_hip_arith.cpp |
Lowers f64 constants and reciprocal. |
platform/cpp/core/lower_hip/core.cpp |
Adds HIP double types and literals. |
platform/cpp/core/ir/ir_tile.cpp |
Supports f64 LDS vectors. |
platform/cpp/core/ir/ir_mem.cpp |
Supports f64 global vectors. |
platform/cpp/core/ir/ir_flow.cpp |
Implements f64 warp shuffling. |
platform/cpp/core/ir/ir_arith.cpp |
Adds f64 constants. |
platform/cpp/core/ir/core_types.cpp |
Registers the F64 singleton and aliases. |
platform/cpp/core/dtypes.cpp |
Adds C++ fp64 dtype metadata. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
6e0880c to
127dd56
Compare
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
The production fusion path silently selects bf16 for f64 inputs, and the public scalar enum change breaks existing ABI tag values.
Review effort: Balanced
Findings: 3
Open (6)
|
The FP64 arithmetic, load/store, cast, and reduce scope looks reasonable. I’m comfortable keeping the following as non-blocking follow-ups, with the current limitations documented explicitly:
These can be tracked separately without expanding this PR, provided unsupported FP64 operations continue to fail explicitly and the support claim reflects the implemented subset. |
2bf3e91 to
f4286f5
Compare
|
The new The starting point is From export PYTHONPATH=python
export ROCKE_LLVM_FLAVOR=llvm20
export ROCKE_LLVM_BIN=/path/to/llvm20/bin
python tools/gen_arch_domain.py
python tools/gen_arch_domain.py --check
python -m pytest tests/core/test_arch_domain_artifact.py -qRepeat with matching LLVM22 and LLVM23 toolchains, then commit the updated Use the full generator invocation here: |
Hi, I was going to generate once all the f64 implementation is completed, but yes (shouldn't be causing regression...), I will use the full generator, fix the coverage test, and commit the updated json files. Thanks for providing the commands! |
f4286f5 to
624f3ae
Compare
tenpercent
left a comment
tenpercent
left a comment
There was a problem hiding this comment.
Thanks! I think we can follow up on the rest
… f64 test coverage
… I/O type via dtype_to_ir_type
|
@tenpercent @minseobshin11 I have some cross-over work in one of my prs and I wanted to do a quick check if I could contribute to this so it can own the f64. Nice work |
624f3ae to
6b8fefd
Compare
cgmillette
left a comment
cgmillette
left a comment
There was a problem hiding this comment.
Thanks for the FP64 work. It's a solid base covering both engines, the HIP lowering and the
elementwise/reduce instances. I compared it with the f64 changes in #13434; #13434 willdrop its own f64 core and build on this PR.
Verified on 6b8fefd
- Byte identity is GREEN on all 76 families at llvm20, llvm22 and llvm23.
- The f64 numeric configs pass on gfx90a.
- f64 global loads and stores select 16-byte
dwordx4. - The f64 warp shuffle is correct.
- LDS aligns f64 buffers to 8 bytes in both engines.
Nothing below blocks the approval. Two small asks before merge; everything else we can pick up in follow-ups, including on our side.
Before merge
1. Fusion lowering silently turns f64 graphs into bf16 kernels.
- For a fusion graph with tensor dtype
float64/f64/fp64,ElementwiseLowerer.can_lowerreturnsok=Trueandcandidates()returns anElementwiseSpecwith dtype'bf16'. The cause is the"f16" if dtype in (...) else "bf16"fallback infusion_lowering.py. FusionLegalizer.check_graphwould reject f64, but I couldn't find a production caller that runs it first.- Since the thread says f64 is excluded from fusion, could
can_lowerreject dtypes it can't map, with an exact message?
2. The max/min/relu NaN contract isn't tested.
- On gfx90a, f64
max/minwith NaN, infinity, -0 and subnormal inputs matchtorch.fmax/torch.fminexactly, but differ fromtorch.maximum/torch.minimum:relu(NaN)returns 0.0;- reduce max/min skip NaNs.
- This is the existing
llvm.maxnumbehaviour f16/bf16/f32 already have, so it isn't a regression. numeric.pycompares against the NaN-propagating references while feeding onlytorch.randn, so it cannot see the difference. Could the references betorch.fmax/torch.fmin?- We'll document the contract and add edge-value inputs in a follow-up.
TheRock Submodule Bump ActivityNewest first
|
…-width and packing contracts, and MMA operand typing (#13434) ## 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. ## 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. `iu8` / `iu4` become spellings of `i8` / `i4`. - `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 #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. ##Fixes: ISSUE ID : AICK-2568, AICK-1713 ## 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 - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com>
…an as a division (#13887) ## Motivation Follow-ups to the base f64 support (#13141), agreed in its review: f64 `abs` took three max instructions where one sign-bit clear does, f64 scalar loads and stores claimed 1-byte alignment, the f64 `mean` rounded twice, two f64 tests could not fail for the reason they named, and the docs' f64 support lists and the LDS pool alignment docstring were out of date. ISSUE ID : AICK-2638 ## Technical Details All changes are in both engines with byte-identical IR. - **fabs / abs:** the `fabs` builder op lowers f64 (`llvm.fabs.f64`; HIP's double `fabs`). The elementwise `abs` uses it for f64 instead of `max(x, -x)`. f32, f16 and bf16 are unchanged. The `fabs.f64` key is measured into the llvm20, llvm22 and llvm23 arch-domain artifacts. A new f64 abs elementwise parity configuration covers the new emission in the byte-identity gate. - **8-byte alignment:** new typed load `global_load_f64` (default align 8); `TensorView.load_scalar` uses it for f64 and `store_scalar` defaults f64 stores to align 8 (the elementwise scalar tail); the reduce2d f64 output store passes align 8. Other dtypes are unchanged. - **mean:** the reduce2d f64 mean divides the sum by N (one rounding) instead of multiplying by 1/N. f16 / bf16 keep the f32 multiply (unchanged IR). - **Tests:** the f64 reduce test now checks that the reduction accumulates in double (combine op, LDS partials, LDS stores), not only that `double` appears; the f64 rejection tests pin the full messages and the `is_valid_spec` reasons. - **Docs:** the f64 rejection lists (`op_vocabulary.md`, `ir_model.md`, `dtype_support.md`) now say tanh is rejected by the builder and exp2_fast has no HIP lowering for any dtype; the LDS pool docstring states the 8-byte f64 alignment, pinned by new tests in both engines. ## Test Plan - Both lowerings of f64 fabs in Python and C++; the f64 abs kernel calls `llvm.fabs.f64` with no `maxnum`. - Every scalar f64 global access of the elementwise tail and the reduce output store claims align 8. - The f64 mean emits one `fdiv double` and no `fmul double`. - GPU on gfx90a: f64 abs, neg, add, mul and copy, including the scalar tail, bit-exact against torch with -0.0, ±inf, NaN, subnormals and the largest double; reduce2d f64 sum / max / mean against torch. ## Test Result - GPU checks pass (bit-exact where exact). - Byte identity green at llvm20, llvm22 and llvm23; ctest passes; the representative-IR golden is unchanged. - Full platform suite: no new failures against develop (the remaining failures also fail on develop and are environment-dependent). - Each commit builds and passes its own tests; pre-commit passes. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com>
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>
ISSUE ID : AICK-1713
Motivation
rocKE has no FP64 device path: there is no
f64IR type, and every family's dtype gate stops at fp16/bf16/fp8/fp32. This PR adds the FP64 core (type, constants, memory, cross-lane, LLVM and HIP lowering) and enables it for theelementwiseandreduce2dfamilies. f64 is computed natively in f64, not promoted to f32.Technical Details
All IR changes are made in both the Python and C++ engines.
F64scalar type withfp64/f64/doubledtype aliases, and aconst_f64builder. f64 literals are emitted as exact IEEE-754 hex in LLVM IR and asdoubleliterals in HIP.f64 -> doublein the LLVM and HIP type maps, plus themaxnum.f64/minnum.f64declarations.vector_add/vector_sub/vector_mullower tofadd/fsub/fmul.math.rcplowers to1.0 / xin f64 on HIP, matching the LLVMfdiv.exp2,exp2_fast,log2,sqrt,rsqrt,rcp_fast,tanh,fabs,fma,fmax3,fmin3andvector_fma. There are no f32↔f64 casts.warp_shuffle_xoraccepts f64 by splitting the value into two i32 halves and shuffling each on the existing 32-bit path.TileWindow.load_vec_scalars, a vector load split into scalars in the storage dtype (the counterpart ofload_vec_as_f32without promotion).sweep_row_chunksand the block LDS reduction keep f64 values in f64.dtype="f64"withvec=2. Allowed ops arecopy,neg,abs,relu,add,sub,mul,maxandmin. The exp2-based activations are rejected byis_valid_spec.dtype="f64"withvec=2. It accumulates in f64, including the LDS scratch, the ±inf identities and themeanscale.dtype_to_ir_type/rocke_dtype_to_ir_type.io_ir_typeinhelpers/io.pyis unchanged (f16/bf16 storage with f32 compute), so f64 does not reach other families through it.intrinsic_arch_domain.llvm{20,22,23}.jsonregenerated withtools/gen_arch_domain.pyon each flavor's toolchain; the only change is the newmaxnum.f64/minnum.f64rows.SUPPORT_MATRIX.md.Test Plan
tools/check_byte_identity.py) on llvm20, llvm22 and llvm23.elementwise_emit(add,relu,max) andreduce_emit(sum,max,mean,min,prod).ctest) and the full pytest suite on gfx942, compared againstdevelop.tests/core/test_arch_domain_artifact.pyon llvm20, llvm22 and llvm23, each with its matching toolchain.rocke_f64_loweringctest:is_valid_specrejects transcendental ops andvec > 2for f64.rcpdivides in double.fadd/fsub/fmul.tests/instances/differential/numeric.py) on gfx942:add,sub,mul,min,relu,abs) and reduce (sum,max,mean,min).rtol = atol = 1e-12(elementwise) and1e-10(reduce).Test Result
Rebased on
develop(1e094cb6b3).develop.maxandmin, and max relative error 1.8e-13 forsumandmean.develop.Next step: the gfx942 FP64 MFMA atoms (
v_mfma_f64_16x16x4_f64,v_mfma_f64_4x4x4_4b_f64), then f64 UniversalGemm.Submission Checklist