[SYCL][HIP] Add gfx90a joint_matrix float/float MFMA combinations#22719
Open
zjin-lcf wants to merge 5 commits into
Open
[SYCL][HIP] Add gfx90a joint_matrix float/float MFMA combinations#22719zjin-lcf wants to merge 5 commits into
zjin-lcf wants to merge 5 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>
CDNA2 (gfx90a / MI2xx) supports the one-block F32 MFMA instructions V_MFMA_F32_16X16X4F32 and V_MFMA_F32_32X32X2F32 (Table 25 of the MI200 CDNA2 ISA), but joint_matrix only exposed fp16/bf16/int8/fp64. Enable the missing fp32 (float in, float out) combinations, shapes 16x16x4 and 32x32x2, mirroring the existing gfx942 support: - matrix-hip.hpp: drop the inner gfx94x-only guard so the fp32 mad path also compiles for gfx90a (the builtins exist on all supported CDNA parts). - static-query-use.hpp: add float/float to the gfx90a type and combination checks. - device_impl.hpp: add the two fp32 combinations to the gfx90a runtime matrix_combinations list. - Update the AMD gfx90a supported-combinations doc table. - Add/extend tests: check_device_code (matrix-hip-float-float-test.cpp), E2E joint_matrix_hip_gfx90a.cpp, runtime_query_hip_gfx90a.cpp and compile-query-hip-gfx90a.cpp. Verified on AMD Instinct MI210 (gfx90a): the E2E and runtime-query tests pass, including numerical correctness of the fp32 MFMA results. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
|
Let's wait for #22682 to be merged before we add more changes on top of that. So review becomes easier. |
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
2 participants
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
V_MFMA_F32_16X16X4F32andV_MFMA_F32_32X32X2F32(Table 25 of the MI200 CDNA2 ISA), butjoint_matrixonly exposed fp16/bf16/int8/fp64.madpath, add the static query + runtimematrix_combinationsentries, and update the docs table.Stacked on
Test plan
matrix-hip-float-float-test.cpp(check_device_code) — FileCheck formfma.f32.16x16x4f32/mfma.f32.32x32x2f32.joint_matrix_hip_gfx90a.cpp,runtime_query_hip_gfx90a.cpp, andcompile-query-hip-gfx90a.cpp.