[SYCL][HIP] Add joint_matrix support for AMD gfx942 (CDNA3 / MI300) by zjin-lcf · Pull Request #22682 · intel/llvm

zjin-lcf · 2026-07-20T03:15:33Z

Summary

Enables the joint_matrix extension on gfx940 / gfx941 / gfx942 (CDNA3, MI300).

matrix-hip.hpp — add the CDNA3 int8 MFMA shapes 16x16x32 and 32x32x16 with packed i64 operands (8 int8 per work-item) and the __builtin_amdgcn_mfma_i32_{16x16x32,32x32x16}_i8 mad 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-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 archs (gfx950, gfx1200, ...) with no further changes.
New gfx942 check_device_code codegen tests, a compile-time query test, and e2e tests (mfma/copy/fill/apply, half, runtime combination query).

Dependencies

Stacked on [SYCL][HIP] Fix AMD joint_matrix use::a row-major load and multi-sub-group lane indexing #22681 (AMD load-path correctness fixes) — review/merge that first; the diff here includes its commit until it merges.
Depends on [SYCL] Fix AMD GPU architecture detection #22670 ("[SYCL] Fix AMD GPU architecture detection") for the runtime combination query to report gfx942 correctly on hardware. The architecture-detection change is intentionally not duplicated here.

Test plan

gfx942 check_device_code codegen tests pass (--param amd_arch=gfx942).
gfx942 e2e tests (mfma/half/runtime-query) pass on MI300A with [SYCL] Fix AMD GPU architecture detection #22670 + [SYCL][HIP] Fix AMD joint_matrix use::a row-major load and multi-sub-group lane indexing #22681 applied.
Existing gfx90a codegen + e2e tests still pass.
CI (codegen tests are hardware-independent; e2e require gfx942 hardware).

…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 · 2026-07-20T14:26:43Z

+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,


What about on gfx940 and gfx941?
they should also be added even if they support the same thing as 942

dkhaldi · 2026-07-20T18:33:20Z

                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>


The test seems to be passing before.
How come there was no issues related to the col vs row major shapes?

"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 · 2026-07-21T14:09:47Z

+}
+
+template <typename Ta, typename Tc>
+constexpr bool are_types_valid_amd_gfx942() {


Can you reuse are_types_valid_amd_gfx90a here? The types are shared between the 4 variants, right?

Thank you for the question/suggestion. I tried to push another PR that add some combinations for gfx90a. These combinations are also supported in gfx942. So the types are shared. For the pending PRs, may we reuse the function after the current PRs are processed ?

Sure but let's wait for this one to be merged then have the new PRs ready to review. Review becomes easier that way.

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>

dkhaldi

LGTM

Co-authored-by: Cursor <cursoragent@cursor.com>

zjin-lcf and others added 2 commits July 19, 2026 20:11

zjin-lcf requested review from a team as code owners July 20, 2026 03:15

zjin-lcf requested review from kweronsx and sergey-semenov July 20, 2026 03:15

dkhaldi reviewed Jul 20, 2026

View reviewed changes

dkhaldi reviewed Jul 21, 2026

View reviewed changes

zjin-lcf mentioned this pull request Jul 22, 2026

[SYCL][HIP] Add gfx90a joint_matrix float/float MFMA combinations #22719

Open

dkhaldi approved these changes Jul 22, 2026

View reviewed changes

[SYCL][HIP] Fix clang-format issues

290d359

Co-authored-by: Cursor <cursoragent@cursor.com>

zjin-lcf requested a review from a team as a code owner July 22, 2026 16:26

[SYCL][HIP] Fix clang-format issues in joint_matrix HIP tests

e6b6918

Co-authored-by: Cursor <cursoragent@cursor.com>

zjin-lcf mentioned this pull request Jul 22, 2026

[SYCL][HIP] Add fp8 (E4M3) / bf8 (E5M2) joint_matrix support on AMD gfx942 #22731

Draft

4 tasks

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[SYCL][HIP] Add joint_matrix support for AMD gfx942 (CDNA3 / MI300)#22682

[SYCL][HIP] Add joint_matrix support for AMD gfx942 (CDNA3 / MI300)#22682
zjin-lcf wants to merge 6 commits into
intel:syclfrom
zjin-lcf:amd-matrix-gfx942

zjin-lcf commented Jul 20, 2026

Uh oh!

dkhaldi Jul 20, 2026

Uh oh!

dkhaldi Jul 20, 2026

Uh oh!

zjin-lcf Jul 20, 2026

Uh oh!

dkhaldi Jul 21, 2026

Uh oh!

zjin-lcf Jul 22, 2026

Uh oh!

dkhaldi Jul 22, 2026 •

edited

Loading

Uh oh!

dkhaldi left a comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

zjin-lcf commented Jul 20, 2026

Summary

Dependencies

Test plan

Uh oh!

dkhaldi Jul 20, 2026

Choose a reason for hiding this comment

Uh oh!

dkhaldi Jul 20, 2026

Choose a reason for hiding this comment

Uh oh!

zjin-lcf Jul 20, 2026

Choose a reason for hiding this comment

Uh oh!

dkhaldi Jul 21, 2026

Choose a reason for hiding this comment

Uh oh!

zjin-lcf Jul 22, 2026

Choose a reason for hiding this comment

Uh oh!

dkhaldi Jul 22, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

dkhaldi left a comment

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

dkhaldi Jul 22, 2026 •

edited

Loading