[SYCL][HIP] Add joint_matrix support for AMD gfx942 (CDNA3 / MI300)#22682
Open
zjin-lcf wants to merge 3 commits into
Open
[SYCL][HIP] Add joint_matrix support for AMD gfx942 (CDNA3 / MI300)#22682zjin-lcf wants to merge 3 commits into
zjin-lcf wants to merge 3 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>
dkhaldi
reviewed
Jul 20, 2026
| template <typename Ta, typename Tb, typename Tc, typename Td, size_t sM, | ||
| size_t sN, size_t sK> | ||
| struct matrix_params< | ||
| architecture::amd_gpu_gfx942, Ta, Tb, Tc, Td, sM, sN, sK, |
Contributor
There was a problem hiding this comment.
What about on gfx940 and gfx941?
they should also be added even if they support the same thing as 942
| joint_matrix<sub_group, InType, use::b, K, N, layout::row_major> | ||
| sub_b; | ||
| joint_matrix<sub_group, InType, use::a, M, K, layout::col_major> | ||
| joint_matrix<sub_group, InType, use::a, M, K, layout::row_major> |
Contributor
There was a problem hiding this comment.
The test seems to be passing before.
How come there was no issues related to the col vs row major shapes?
Contributor
Author
There was a problem hiding this comment.
"The old col_major load was latently wrong but undetectable: in this test A is filled entirely with 1s, so a col-major vs row-major load of A reads identical values, masking the layout mistake. The host reference computes with row-major indexing (A[m*K+k]), and the companion joint_matrix_hip_mfma.hpp already uses layout::row_major for use::a, so switching to row_major makes the layout correct and consistent — the test passed before only because of the uniform input data"
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>
dkhaldi
reviewed
Jul 21, 2026
| } | ||
|
|
||
| template <typename Ta, typename Tc> | ||
| constexpr bool are_types_valid_amd_gfx942() { |
Contributor
There was a problem hiding this comment.
Can you reuse are_types_valid_amd_gfx90a here? The types are shared between the 4 variants, right?
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
Enables the
joint_matrixextension on gfx940 / gfx941 / gfx942 (CDNA3, MI300).matrix-hip.hpp— add the CDNA3 int8 MFMA shapes16x16x32and32x32x16with packed i64 operands (8 int8 per work-item) and the__builtin_amdgcn_mfma_i32_{16x16x32,32x32x16}_i8mad paths; 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-timematrix_paramsspecializations.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 theamd_archlit param (defaultgfx90a) and expose ahip-arch-<arch>feature so arch-specific codegen tests gate themselves. Existing gfx90a codegen tests nowREQUIRE hip-arch-gfx90a. This keeps the harness scalable to future archs (gfx950, gfx1200, ...) with no further changes.check_device_codecodegen tests, a compile-time query test, and e2e tests (mfma/copy/fill/apply, half, runtime combination query).Dependencies
Test plan
check_device_codecodegen tests pass (--param amd_arch=gfx942).