[SYCL][HIP] Add fp8 (E4M3) / bf8 (E5M2) joint_matrix support on AMD gfx942#22731
Draft
zjin-lcf wants to merge 10 commits into
Draft
[SYCL][HIP] Add fp8 (E4M3) / bf8 (E5M2) joint_matrix support on AMD gfx942#22731zjin-lcf wants to merge 10 commits into
zjin-lcf wants to merge 10 commits into
Conversation
…group lane indexing Two correctness fixes in the AMD (HIP) joint_matrix backend load path: 1. The use::a row-major load loaded the A multiplicand transposed. The access pattern was chosen from the layout enum alone, but "row-major" implies opposite in-memory stride patterns for A and B (the free dimension is strided for A, contiguous for B). Select the pattern from the combined (use, layout) condition so both layouts are correct for A. The double and non-double paths are unified (Size == 1 reduces to the former double case), which also fixes use::b col-major for double. 2. The per-lane fragment index was computed from the work-group linear id (get_group_linear_id() * sub_group_size + lane) instead of the sub-group local id. With a single sub-group per work-group this is accidentally correct, but kernels launching multiple sub-groups per work-group computed out-of-range element offsets, giving wrong results or memory faults. The hip Inputs helpers previously declared use::a as col_major while storing A row-major, relying on the old (buggy) behavior; they now use row_major. Adds two regression e2e tests: all four A/B layout combinations, and a tiled GEMM with multiple sub-groups per work-group. Co-authored-by: Cursor <cursoragent@cursor.com>
Enable the joint_matrix extension on gfx940/gfx941/gfx942 (CDNA3):
- matrix-hip.hpp: add the CDNA3 int8 MFMA shapes 16x16x32 and 32x32x16 with
packed i64 operands (8 int8 per work-item) and the corresponding
__builtin_amdgcn_mfma_i32_{16x16x32,32x32x16}_i8 mad paths; the fp16/bf16/
fp64 shapes are shared with gfx90a.
- matrix-unified.hpp: extend the backend include guard to CDNA3.
- static-query-use.hpp: add the compile-time matrix_params specializations.
- device_impl.hpp: add the runtime combination list for gfx940/941/942.
- test/lit.cfg.py: make the AMD codegen suite target-arch configurable via the
`amd_arch` lit param (default gfx90a) and expose a `hip-arch-<arch>` feature
so arch-specific codegen tests gate themselves; existing gfx90a codegen tests
now REQUIRE hip-arch-gfx90a. This keeps the harness scalable to future
architectures with no further changes.
- Add gfx942 check_device_code codegen tests, a compile-time query test, and
e2e tests (mfma/copy/fill/apply, half, and the runtime combination query).
Depends on intel#22670 (AMD GPU architecture detection) so the runtime combination
query reports gfx942 correctly on hardware.
Co-authored-by: Cursor <cursoragent@cursor.com>
Address review feedback: gfx940 and gfx941 are CDNA3 parts with the same joint_matrix support as gfx942, but only gfx942 had matrix_params specializations. Add the default-values and validation matrix_params specializations for amd_gpu_gfx940 and amd_gpu_gfx941 (reusing the gfx942 type/combination helpers), extend the compile-time query test to cover all three architectures, and gate the CDNA3 check_device_code tests on any of gfx940/gfx941/gfx942 since they emit identical codegen. Co-authored-by: Cursor <cursoragent@cursor.com>
Add the single-block F32 MFMA shapes from CDNA3 ISA Table 28 that were missing from the gfx942 joint_matrix support: 16x16x4 and 32x32x2 with float A/B and float accumulator. - static-query-use.hpp: accept float/float (16x16x4, 32x32x2) in is_combination_valid_amd_gfx942, add float to are_types_valid_amd_gfx942, and default K to 4 for float (gfx942/940/941). - matrix-hip.hpp: add float a/b fragment overloads and lower to __builtin_amdgcn_mfma_f32_16x16x4f32 / _32x32x2f32. - Tests: extend compile-query-hip-gfx942.cpp and add matrix-hip-gfx942-float-float-test.cpp codegen check. Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Add joint_matrix support for the 8-bit floating point MFMA instructions on AMD CDNA3 (gfx942/gfx940/gfx941). Uses the experimental fp8_e4m3 (E4M3) and fp8_e5m2 (E5M2) types directly as joint_matrix element types with an fp32 accumulator, for the 16x16x32 and 32x32x16 shapes. The A and B operand formats are independent, so all four (E4M3, E5M2) pairs are supported per shape (8 combinations). - matrix-hip.hpp: fp8 fragment overloads and an fp8/bf8 MFMA branch in joint_matrix_mad_hip with separate A/B multiplicand types. - matrix-unified.hpp: relax the HIP Ta == Tb gate to allow mixed fp8 pairs. - query-types.hpp: matrix_type::fp8_e4m3 / fp8_e5m2 plus string mappings. - static-query-use.hpp: extend gfx942/940/941 matrix_params (incl. mixed A/B). - device_impl.hpp: add the 8 fp8 runtime combinations. - Tests: e2e (8 combos), check_device_code FileCheck, and compile-query. Co-authored-by: Cursor <cursoragent@cursor.com>
checkDevSupportJointMatrix compares the kernel's joint_matrix property strings against the device's matrix_combinations. Without recognizing the fp8 type strings the runtime rejects fp8 joint_matrix kernels even on gfx942. Add the missing cases to convertMatrixTypeStringMatrixTypeEnumValue. Co-authored-by: Cursor <cursoragent@cursor.com>
The joint_matrix fp8 element types are the OCP E4M3/E5M2 formats, but the CDNA3 (gfx940/gfx941/gfx942) fp8/bf8 MFMA instructions read their operands as fnuz, which uses a +1 exponent bias. Re-encode the packed i64 operands in joint_matrix_mad_hip with a branchless, integer-only, fully-unrolled transform (no float round-trip, no lookup table): subnormals shift left by one, normals bump the exponent field, and zero/NaN/inf map to the fnuz conventions. This fixes the numerical mismatch observed on MI300A. Co-authored-by: Cursor <cursoragent@cursor.com>
Run every fp8/bf8 joint_matrix combination with both a row_major and a col_major accumulator store to cover the transposed store path. Co-authored-by: Cursor <cursoragent@cursor.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds SYCL
joint_matrixsupport for the experimental 8-bit floating-point types (fp8_e4m3/fp8_e5m2fromsycl::ext::oneapi::experimental) on AMD gfx942 (CDNA3 / MI300) matrix cores, using thef32 = fp8/bf8 x fp8/bf8MFMA instructions for the16x16x32and32x32x16shapes with anf32accumulator. All eight A/B format pairings (E4M3/E5M2 x E4M3/E5M2) are supported, including mixed formats.Key points:
static-query) and runtime (matrix_combinations) query paths and the device-codematrix_typestring parsing for the new fp8 types.Note
This PR is stacked on #22682 and its first four commits belong to that PR; they will drop out of this diff once #22682 is merged. Review the top four commits (fp8 support) here.
Test plan
sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-fp8-test.cpp- FileCheck verifies all 8 fp8/bf8 MFMA intrinsics are emitted.sycl/test/matrix/hip/compile-query-hip-gfx942.cpp- compile-time static-query assertions for fp8 combinations.sycl/test-e2e/Matrix/joint_matrix_hip_fp8_gfx942.cpp- e2e over all 8 combos x 2 shapes x KX 1-4 x {row_major, col_major} store; passes on MI300A.