From 6eafc610ea5b49fa4aa81e3003e3b1fb7d969706 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sun, 19 Jul 2026 20:11:18 -0700 Subject: [PATCH 1/5] [SYCL][HIP] Fix AMD joint_matrix use::a row-major load and multi-sub-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 --- .../sycl/ext/oneapi/matrix/matrix-hip.hpp | 83 +++++++++----- .../Matrix/Inputs/joint_matrix_hip_apply.hpp | 2 +- .../Matrix/Inputs/joint_matrix_hip_copy.hpp | 2 +- .../Matrix/Inputs/joint_matrix_hip_fill.hpp | 2 +- .../Matrix/Inputs/joint_matrix_hip_mfma.hpp | 2 +- .../Matrix/joint_matrix_hip_all_layouts.cpp | 106 ++++++++++++++++++ .../joint_matrix_hip_multi_subgroup.cpp | 99 ++++++++++++++++ 7 files changed, 263 insertions(+), 33 deletions(-) create mode 100644 sycl/test-e2e/Matrix/joint_matrix_hip_all_layouts.cpp create mode 100644 sycl/test-e2e/Matrix/joint_matrix_hip_multi_subgroup.cpp diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp index 6f55ae3f1d6cf..de70895bf3eb0 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp @@ -117,8 +117,11 @@ void load_accumulator_layoutT( S, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, sycl::ext::oneapi::experimental::matrix::layout::dynamic> &res, multi_ptr src, size_t stride, Group &sg) { - const auto idx = sg.get_group_linear_id() * sg.get_local_range()[0] + - sg.get_local_linear_id(); + // The fragment lane index is the work-item's position within its own + // sub-group (0 .. sub_group_size-1). It must NOT include the sub-group's + // index within the work-group, otherwise kernels that launch more than one + // sub-group per work-group compute out-of-range element offsets. + const auto idx = sg.get_local_linear_id(); if constexpr (std::is_same_v) { const auto thread_x = idx % N; @@ -211,33 +214,52 @@ template < void load_multiplicand_hip(joint_matrix_hip &res, multi_ptr src, size_t stride, Group &sg) { - const auto idx = sg.get_group_linear_id() * sg.get_local_range()[0] + - sg.get_local_linear_id(); - - if constexpr (std::is_same_v) { - if constexpr (Layout == - sycl::ext::oneapi::experimental::matrix::layout::row_major) { - res.wi_marray[0] = src[idx]; - } else { - res.wi_marray[0] = src[(idx % M) * stride + idx / M]; + // The fragment lane index is the work-item's position within its own + // sub-group (0 .. sub_group_size-1). It must NOT include the sub-group's + // index within the work-group, otherwise kernels that launch more than one + // sub-group per work-group compute out-of-range element offsets. + const auto idx = sg.get_local_linear_id(); + + // The AMD MFMA fragment layout is described in terms of the matrix's "free" + // dimension (rows M for the A multiplicand, columns N for the B multiplicand) + // and the shared contraction dimension K. Whether the free dimension is the + // strided one in memory depends on BOTH the use and the layout: + // - A row_major: element (m,k) at m*stride + k -> free dim (m) is strided. + // - A col_major: element (m,k) at k*stride + m -> free dim (m) is contiguous. + // - B row_major: element (k,n) at k*stride + n -> free dim (n) is contiguous. + // - B col_major: element (k,n) at n*stride + k -> free dim (n) is strided. + // So the strided-free-dim access pattern applies to (A, row_major) and + // (B, col_major); the contiguous-free-dim pattern applies to (A, col_major) + // and (B, row_major). Selecting the pattern from this combined condition + // (rather than from the layout alone) makes both layouts behave correctly for + // the A multiplicand, which previously loaded A transposed for row_major. + constexpr bool StridedFreeDim = + (Use == sycl::ext::oneapi::experimental::matrix::use::a) + ? (Layout == + sycl::ext::oneapi::experimental::matrix::layout::row_major) + : (Layout == + sycl::ext::oneapi::experimental::matrix::layout::col_major); + + // Free dimension: M for the A multiplicand, N for the B multiplicand. The + // wavefront is split into (WAVEFRONT_SIZE / FreeDim) groups along the + // contraction dimension K, and each work-item holds one contiguous K-block of + // `Size` elements (Size == 1 for the double shapes). + constexpr int FreeDim = + (Use == sycl::ext::oneapi::experimental::matrix::use::a) ? M : N; + constexpr int Size = (M * N) / WAVEFRONT_SIZE; + + const auto thread_x = idx % FreeDim; + const auto thread_y = idx / FreeDim; + + if constexpr (StridedFreeDim) { + for (int i = 0; i < Size; ++i) { + const int c_idx = thread_x * stride + i + thread_y * Size; + res.wi_marray[i] = src[c_idx]; } } else { - constexpr int Dim = (M == 16) ? 16 : 32; - - const auto thread_x = idx % Dim; - const auto thread_y = idx / Dim; - - if constexpr (Layout == - sycl::ext::oneapi::experimental::matrix::layout::col_major) { - for (int i = 0; i < 4; ++i) { - const int c_idx = thread_x * stride + i + thread_y * 4; - res.wi_marray[i] = src[c_idx]; - } - } else { - for (int i = 0; i < 4; ++i) { - const int r_idx = thread_x + i * stride + thread_y * stride * 4; - res.wi_marray[i] = src[r_idx]; - } + for (int i = 0; i < Size; ++i) { + const int r_idx = thread_x + i * stride + thread_y * stride * Size; + res.wi_marray[i] = src[r_idx]; } } } @@ -251,8 +273,11 @@ void store_layoutT( T, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, sycl::ext::oneapi::experimental::matrix::layout::dynamic> &src, multi_ptr dst, size_t stride, Group &sg) { - const auto idx = sg.get_group_linear_id() * sg.get_local_range()[0] + - sg.get_local_linear_id(); + // The fragment lane index is the work-item's position within its own + // sub-group (0 .. sub_group_size-1). It must NOT include the sub-group's + // index within the work-group, otherwise kernels that launch more than one + // sub-group per work-group compute out-of-range element offsets. + const auto idx = sg.get_local_linear_id(); if constexpr (std::is_same_v) { const auto thread_x = idx % N; diff --git a/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_apply.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_apply.hpp index 01f187f2368a7..307964fc0c00e 100644 --- a/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_apply.hpp +++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_apply.hpp @@ -60,7 +60,7 @@ void hip_matrix_apply() { joint_matrix sub_c; joint_matrix sub_b; - joint_matrix + joint_matrix sub_a; joint_matrix_load( diff --git a/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_copy.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_copy.hpp index e56fecf8df5b3..a560d3400dd3b 100644 --- a/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_copy.hpp +++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_copy.hpp @@ -69,7 +69,7 @@ void hip_matrix_copy() { sub_c_copy; joint_matrix sub_b, sub_b_copy; - joint_matrix + joint_matrix sub_a, sub_a_copy; joint_matrix_load( diff --git a/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_fill.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_fill.hpp index f997836670ebf..bca0909819cc2 100644 --- a/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_fill.hpp +++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_fill.hpp @@ -61,7 +61,7 @@ void hip_matrix_fill() { sub_c{}; joint_matrix sub_b{}; - joint_matrix + joint_matrix sub_a{}; joint_matrix_fill(sg, sub_a, 1); diff --git a/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_mfma.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_mfma.hpp index d066350b638f0..9d46c271037e0 100644 --- a/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_mfma.hpp +++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_hip_mfma.hpp @@ -68,7 +68,7 @@ void hip_matrix_mfma() { joint_matrix sub_c; joint_matrix sub_b; - joint_matrix + joint_matrix sub_a; joint_matrix_load( diff --git a/sycl/test-e2e/Matrix/joint_matrix_hip_all_layouts.cpp b/sycl/test-e2e/Matrix/joint_matrix_hip_all_layouts.cpp new file mode 100644 index 0000000000000..512d7de3f70df --- /dev/null +++ b/sycl/test-e2e/Matrix/joint_matrix_hip_all_layouts.cpp @@ -0,0 +1,106 @@ +//===---joint_matrix_hip_all_layouts.cpp - DPC++ joint_matrix-------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// RUN: %clangxx -fsycl -fsycl-targets=amdgcn-amd-amdhsa %amd_arch_options %s -o %t.out +// RUN: %{run} %t.out + +// REQUIRES: target-amd + +// Regression test: a 16x16x16 bf16 -> float tile computed for every +// combination of A/B memory layouts (row/col major), each stored in memory +// according to its declared layout, and checked against a standard row-major +// C = A*B reference. Guards against the AMD `use::a` load-path layout bug where +// the row-major A multiplicand was loaded transposed. + +#include +#include +#include + +#include +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; +using sycl::ext::oneapi::bfloat16; + +constexpr int M = 16, N = 16, K = 16; + +template bool run_combo() { + std::vector Alog(M * K), Blog(K * N), Ref(M * N, 0.f); + for (int i = 0; i < M * K; ++i) + Alog[i] = static_cast((i * 3 + 1) % 7) - 3.f; + for (int i = 0; i < K * N; ++i) + Blog[i] = static_cast((i * 2 + 1) % 5) - 2.f; + for (int m = 0; m < M; ++m) + for (int n = 0; n < N; ++n) { + float e = 0.f; + for (int k = 0; k < K; ++k) + e += Alog[m * K + k] * Blog[k * N + n]; + Ref[m * N + n] = e; + } + + const int lda = (AL == layout::row_major) ? K : M; + const int ldb = (BL == layout::row_major) ? N : K; + std::vector Amem(M * K), Bmem(K * N); + for (int m = 0; m < M; ++m) + for (int k = 0; k < K; ++k) + Amem[(AL == layout::row_major) ? m * lda + k : k * lda + m] = + bfloat16(Alog[m * K + k]); + for (int k = 0; k < K; ++k) + for (int n = 0; n < N; ++n) + Bmem[(BL == layout::row_major) ? k * ldb + n : n * ldb + k] = + bfloat16(Blog[k * N + n]); + std::vector C(M * N, 0.f); + + { + queue q; + buffer bA(Amem.data(), range{M * K}); + buffer bB(Bmem.data(), range{K * N}); + buffer bC(C.data(), range{M * N}); + q.submit([&](handler &h) { + accessor accA{bA, h, read_only}; + accessor accB{bB, h, read_only}; + accessor accC{bC, h, write_only}; + h.parallel_for( + nd_range<2>{{4, 16}, {4, 16}}, [=](nd_item<2> it) { + auto sg = it.get_sub_group(); + joint_matrix c; + joint_matrix a; + joint_matrix b; + joint_matrix_fill(sg, c, 0.f); + joint_matrix_load( + sg, a, accA.template get_multi_ptr(), + lda); + joint_matrix_load( + sg, b, accB.template get_multi_ptr(), + ldb); + joint_matrix_mad(sg, c, a, b, c); + joint_matrix_store( + sg, c, accC.template get_multi_ptr(), N, + layout::row_major); + }); + }).wait(); + } + + for (int i = 0; i < M * N; ++i) + if (std::fabs(C[i] - Ref[i]) > 2.f) + return false; + return true; +} + +int main() { + bool rr = run_combo(); + bool rc = run_combo(); + bool cr = run_combo(); + bool cc = run_combo(); + assert(rr && "A=row B=row"); + assert(rc && "A=row B=col"); + assert(cr && "A=col B=row"); + assert(cc && "A=col B=col"); + return 0; +} diff --git a/sycl/test-e2e/Matrix/joint_matrix_hip_multi_subgroup.cpp b/sycl/test-e2e/Matrix/joint_matrix_hip_multi_subgroup.cpp new file mode 100644 index 0000000000000..7d76dee85fa51 --- /dev/null +++ b/sycl/test-e2e/Matrix/joint_matrix_hip_multi_subgroup.cpp @@ -0,0 +1,99 @@ +//===---joint_matrix_hip_multi_subgroup.cpp - DPC++ joint_matrix----------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// RUN: %clangxx -fsycl -fsycl-targets=amdgcn-amd-amdhsa %amd_arch_options %s -o %t.out +// RUN: %{run} %t.out + +// REQUIRES: target-amd + +// Regression test: a tiled GEMM whose work-group contains MORE THAN ONE +// sub-group (a 64x64 output tile computed by a 4x4 grid of wavefronts). Guards +// against the AMD fragment lane-index bug where the per-lane offset was derived +// from the work-group linear id instead of the sub-group local id, which made +// every sub-group beyond the first read/write out of range. + +#include +#include +#include + +#include +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; +using sycl::ext::oneapi::bfloat16; + +constexpr int T = 16; // MFMA tile M=N=K +constexpr int TILE = 64; // work-group output tile (TILE x TILE) +constexpr int S = 128; // full square matrix dimension (multiple of TILE) + +int main() { + std::vector Alog(S * S), Blog(S * S), Ref(S * S, 0.f); + for (int i = 0; i < S * S; ++i) + Alog[i] = static_cast((i * 3 + 1) % 7) - 3.f; + for (int i = 0; i < S * S; ++i) + Blog[i] = static_cast((i * 2 + 1) % 5) - 2.f; + for (int m = 0; m < S; ++m) + for (int n = 0; n < S; ++n) { + float e = 0.f; + for (int k = 0; k < S; ++k) + e += Alog[m * S + k] * Blog[k * S + n]; + Ref[m * S + n] = e; + } + + std::vector Amem(S * S), Bmem(S * S); // A row-major, B col-major + for (int m = 0; m < S; ++m) + for (int k = 0; k < S; ++k) + Amem[m * S + k] = bfloat16(Alog[m * S + k]); + for (int k = 0; k < S; ++k) + for (int n = 0; n < S; ++n) + Bmem[n * S + k] = bfloat16(Blog[k * S + n]); + std::vector C(S * S, 0.f); + + { + queue q; + buffer bA(Amem.data(), range{(size_t)S * S}); + buffer bB(Bmem.data(), range{(size_t)S * S}); + buffer bC(C.data(), range{(size_t)S * S}); + // Work-group: {1, TILE/T, (TILE/T)*WAVE} where WAVE=64 -> 16 sub-groups. + const int nwarp = TILE / T; // 4 + range<2> lws{(size_t)nwarp, (size_t)nwarp * 64}; + range<2> gws{(size_t)(S / TILE) * nwarp, (size_t)(S / TILE) * nwarp * 64}; + q.submit([&](handler &h) { + accessor accA{bA, h, read_only}; + accessor accB{bB, h, read_only}; + accessor accC{bC, h, write_only}; + h.parallel_for( + nd_range<2>{gws, lws}, [=](nd_item<2> it) { + auto sg = it.get_sub_group(); + const int warpM = it.get_local_id(0); + const int warpN = it.get_local_id(1) / 64; + const int row = (it.get_group(0) * (TILE / T) + warpM) * T; + const int col = (it.get_group(1) * (TILE / T) + warpN) * T; + auto pa = accA.template get_multi_ptr(); + auto pb = accB.template get_multi_ptr(); + auto pc = accC.template get_multi_ptr(); + joint_matrix c; + joint_matrix a; + joint_matrix b; + joint_matrix_fill(sg, c, 0.f); + for (int kk = 0; kk < S; kk += T) { + joint_matrix_load(sg, a, pa + row * S + kk, S); + joint_matrix_load(sg, b, pb + col * S + kk, S); + joint_matrix_mad(sg, c, a, b, c); + } + joint_matrix_store(sg, c, pc + row * S + col, S, + layout::row_major); + }); + }).wait(); + } + + for (int i = 0; i < S * S; ++i) + assert(std::fabs(C[i] - Ref[i]) < 4.f && "multi-subgroup GEMM mismatch"); + return 0; +} From d7005877329e9c66e30ee615879913f587c8b815 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sun, 19 Jul 2026 20:14:40 -0700 Subject: [PATCH 2/5] [SYCL][HIP] Add joint_matrix support for AMD gfx942 (CDNA3 / MI300) 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-` 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 #22670 (AMD GPU architecture detection) so the runtime combination query reports gfx942 correctly on hardware. Co-authored-by: Cursor --- .../sycl/ext/oneapi/matrix/matrix-hip.hpp | 32 ++++++- .../sycl/ext/oneapi/matrix/matrix-unified.hpp | 3 +- .../ext/oneapi/matrix/static-query-use.hpp | 90 +++++++++++++++++++ sycl/source/detail/device_impl.hpp | 19 ++++ .../Matrix/joint_matrix_hip_gfx942.cpp | 63 +++++++++++++ .../Matrix/joint_matrix_hip_half_gfx942.cpp | 44 +++++++++ .../Matrix/runtime_query_hip_gfx942.cpp | 68 ++++++++++++++ .../matrix/matrix-hip-bfloat16-float-test.cpp | 1 + .../matrix/matrix-hip-double-double-test.cpp | 1 + .../matrix-hip-gfx942-bfloat16-float-test.cpp | 65 ++++++++++++++ .../matrix-hip-gfx942-double-double-test.cpp | 37 ++++++++ .../matrix-hip-gfx942-half-float-test.cpp | 64 +++++++++++++ .../matrix-hip-gfx942-int8-int32-test.cpp | 65 ++++++++++++++ .../hip/matrix/matrix-hip-half-float-test.cpp | 1 + .../hip/matrix/matrix-hip-int8-int32-test.cpp | 1 + sycl/test/lit.cfg.py | 26 ++++-- .../matrix/hip/compile-query-hip-gfx942.cpp | 28 ++++++ 17 files changed, 600 insertions(+), 8 deletions(-) create mode 100644 sycl/test-e2e/Matrix/joint_matrix_hip_gfx942.cpp create mode 100644 sycl/test-e2e/Matrix/joint_matrix_hip_half_gfx942.cpp create mode 100644 sycl/test-e2e/Matrix/runtime_query_hip_gfx942.cpp create mode 100644 sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-bfloat16-float-test.cpp create mode 100644 sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-double-double-test.cpp create mode 100644 sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-half-float-test.cpp create mode 100644 sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-int8-int32-test.cpp create mode 100644 sycl/test/matrix/hip/compile-query-hip-gfx942.cpp diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp index de70895bf3eb0..e66ec2ba73651 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp @@ -91,6 +91,13 @@ __SYCL_JOINT_MATRIX_OVERLOAD_ARR(int8_t, b, 8, 32, 4) __SYCL_JOINT_MATRIX_OVERLOAD_ARR(int8_t, a, 16, 16, 4) __SYCL_JOINT_MATRIX_OVERLOAD_ARR(int8_t, b, 16, 16, 4) +// gfx942 (CDNA3) int8 MFMA uses larger K dimensions with packed operands: +// 16x16x32 and 32x32x16. Each work-item holds 8 int8 values (one i64). +__SYCL_JOINT_MATRIX_OVERLOAD_ARR(int8_t, a, 16, 32, 8) +__SYCL_JOINT_MATRIX_OVERLOAD_ARR(int8_t, b, 32, 16, 8) +__SYCL_JOINT_MATRIX_OVERLOAD_ARR(int8_t, a, 32, 16, 8) +__SYCL_JOINT_MATRIX_OVERLOAD_ARR(int8_t, b, 16, 32, 8) + #undef __SYCL_JOINT_MATRIX_OVERLOAD_ARR #define __SYCL_JOINT_MATRIX_OVERLOAD_ARR_ACC(TYPE, M, N) \ @@ -370,7 +377,8 @@ void joint_matrix_mad_hip( const joint_matrix_hip< Tc, sycl::ext::oneapi::experimental::matrix::use::accumulator, M, N, sycl::ext::oneapi::experimental::matrix::layout::dynamic> &C) { -#ifdef __gfx90a__ +#if defined(__gfx90a__) || defined(__gfx940__) || defined(__gfx941__) || \ + defined(__gfx942__) if constexpr (std::is_same_v) { if constexpr (M == 16 && N == 16) { auto result = __builtin_amdgcn_mfma_f32_16x16x16f16( @@ -407,6 +415,25 @@ void joint_matrix_mad_hip( std::memcpy(&D.wi_marray, &result, 4 * sizeof(double)); } } else if constexpr (std::is_same_v) { +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + // CDNA3 (gfx942) int8 MFMA: 16x16x32 / 32x32x16 with packed i64 operands + // (8 int8 values per work-item). + if constexpr (M == 16 && N == 16) { + auto result = __builtin_amdgcn_mfma_i32_16x16x32_i8( + *reinterpret_cast(&A.wi_marray), + *reinterpret_cast(&B.wi_marray), + *reinterpret_cast(&C.wi_marray), 0, 0, 0); + std::memcpy(&D.wi_marray, &result, 4 * sizeof(int32_t)); + } else if constexpr (M == 32 && N == 32) { + auto result = __builtin_amdgcn_mfma_i32_32x32x16_i8( + *reinterpret_cast(&A.wi_marray), + *reinterpret_cast(&B.wi_marray), + *reinterpret_cast(&C.wi_marray), 0, 0, 0); + std::memcpy(&D.wi_marray, &result, 16 * sizeof(int32_t)); + } +#else + // CDNA2 (gfx90a) int8 MFMA: 16x16x16 / 32x32x8 with i32 operands + // (4 int8 values per work-item). if constexpr (M == 16 && N == 16) { auto result = __builtin_amdgcn_mfma_i32_16x16x16i8( *reinterpret_cast(&A.wi_marray), @@ -420,8 +447,9 @@ void joint_matrix_mad_hip( *reinterpret_cast(&C.wi_marray), 0, 0, 0); std::memcpy(&D.wi_marray, &result, 16 * sizeof(int32_t)); } +#endif } -#endif // __gfx90a__ +#endif // __gfx90a__ || __gfx942__ } } // namespace detail diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp index f107aabf39115..2a6fcd4e89fa0 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp @@ -13,7 +13,8 @@ #if defined(__SYCL_DEVICE_ONLY__) #if defined(__NVPTX__) #include "matrix-tensorcores.hpp" -#elif defined(__gfx90a__) +#elif defined(__gfx90a__) || defined(__gfx940__) || defined(__gfx941__) || \ + defined(__gfx942__) #include "matrix-hip.hpp" #endif // defined(__NVPTX__) #endif // defined(__SYCL_DEVICE_ONLY__) diff --git a/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp b/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp index 7e0183ae50895..0c1ea1d2c536a 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp @@ -574,6 +574,96 @@ struct matrix_params< using joint_matrix_d = joint_matrix; }; +////////////////////////////////////////////// +/// AMD Matrix Cores - GFX942 architecture /// +////////////////////////////////////////////// + +template +constexpr bool is_combination_valid_amd_gfx942(size_t sM, size_t sN, + size_t sK) { + return (std::is_same_v && std::is_same_v && + ((sM == 32 && sN == 32 && sK == 8) || + (sM == 16 && sN == 16 && sK == 16))) || + (std::is_same_v && std::is_same_v && + ((sM == 32 && sN == 32 && sK == 16) || + (sM == 16 && sN == 16 && sK == 32))) || + (std::is_same_v && std::is_same_v && + ((sM == 32 && sN == 32 && sK == 8) || + (sM == 16 && sN == 16 && sK == 16))) || + (std::is_same_v && std::is_same_v && + (sM == 16 && sN == 16 && sK == 4)); +} + +template +constexpr bool are_types_valid_amd_gfx942() { + return (std::is_same_v && std::is_same_v) || + (std::is_same_v && std::is_same_v) || + (std::is_same_v && std::is_same_v) || + (std::is_same_v && std::is_same_v); +} + +// Default-values query: +// Specialization for when only types are given, need to query only sizes +template +struct matrix_params< + architecture::amd_gpu_gfx942, Ta, Tb, Tc, Td, 0, 0, 0, + typename std::enable_if_t<( + !std::is_same_v && !std::is_same_v && + !std::is_same_v && !std::is_same_v && + std::is_same_v && std::is_same_v)>> { + static_assert( + are_types_valid_amd_gfx942(), + "Invalid types for AMD gfx942, supported types are half, float, " + "int8_t, int32_t, double and bfloat16 "); + + // Default sizes for AMD gfx942 were chosen to represent a square matrix + static constexpr std::size_t M = 16; + static constexpr std::size_t N = 16; + static constexpr std::size_t K = std::is_same_v ? 4 + : std::is_same_v ? 32 + : 16; + + template + using joint_matrix_a = joint_matrix; + template + using joint_matrix_b = joint_matrix; + template + using joint_matrix_c = joint_matrix; + template + using joint_matrix_d = joint_matrix; +}; + +// Validation query +// Specialization when both types and sizes are given +template +struct matrix_params< + architecture::amd_gpu_gfx942, Ta, Tb, Tc, Td, sM, sN, sK, + typename std::enable_if_t<( + !std::is_same_v && !std::is_same_v && + !std::is_same_v && !std::is_same_v && + std::is_same_v && std::is_same_v && sM != 0 && + sN != 0 && sK != 0)>> { + static_assert( + is_combination_valid_amd_gfx942(sM, sN, sK), + "Invalid parameters for AMD gfx942, query valid combinations " + "using: " + "q.get_device().get_info()"); + + static constexpr std::size_t M = sM; + static constexpr std::size_t N = sN; + static constexpr std::size_t K = sK; + + template + using joint_matrix_a = joint_matrix; + template + using joint_matrix_b = joint_matrix; + template + using joint_matrix_c = joint_matrix; + template + using joint_matrix_d = joint_matrix; +}; + ///////////////////////////////////////////////// /// CUDA Tensor Cores - sm70, sm72 and sm80 /// ///////////////////////////////////////////////// diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index 3715e0c0ae1ce..95c484d514042 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -2222,6 +2222,25 @@ class device_impl { {0, 0, 0, 16, 16, 4, matrix_type::fp64, matrix_type::fp64, matrix_type::fp64, matrix_type::fp64}, }; + else if ((architecture::amd_gpu_gfx940 == DeviceArch) || + (architecture::amd_gpu_gfx941 == DeviceArch) || + (architecture::amd_gpu_gfx942 == DeviceArch)) + return { + {0, 0, 0, 32, 32, 8, matrix_type::fp16, matrix_type::fp16, + matrix_type::fp32, matrix_type::fp32}, + {0, 0, 0, 16, 16, 16, matrix_type::fp16, matrix_type::fp16, + matrix_type::fp32, matrix_type::fp32}, + {0, 0, 0, 32, 32, 16, matrix_type::sint8, matrix_type::sint8, + matrix_type::sint32, matrix_type::sint32}, + {0, 0, 0, 16, 16, 32, matrix_type::sint8, matrix_type::sint8, + matrix_type::sint32, matrix_type::sint32}, + {0, 0, 0, 32, 32, 8, matrix_type::bf16, matrix_type::bf16, + matrix_type::fp32, matrix_type::fp32}, + {0, 0, 0, 16, 16, 16, matrix_type::bf16, matrix_type::bf16, + matrix_type::fp32, matrix_type::fp32}, + {0, 0, 0, 16, 16, 4, matrix_type::fp64, matrix_type::fp64, + matrix_type::fp64, matrix_type::fp64}, + }; else if (backend::ext_oneapi_cuda == CurrentBackend) { // TODO: Tho following can be simplified when comparison of // architectures using < and > will be implemented diff --git a/sycl/test-e2e/Matrix/joint_matrix_hip_gfx942.cpp b/sycl/test-e2e/Matrix/joint_matrix_hip_gfx942.cpp new file mode 100644 index 0000000000000..cc46edd104a60 --- /dev/null +++ b/sycl/test-e2e/Matrix/joint_matrix_hip_gfx942.cpp @@ -0,0 +1,63 @@ +//===---joint_matrix_hip_gfx942.cpp - DPC++ joint_matrix-------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// RUN: %clangxx -fsycl -fsycl-targets=amd_gpu_gfx942 %s -o %t.out +// RUN: %{run} %t.out + +// REQUIRES: target-amd +// REQUIRES: arch-amd_gpu_gfx942 + +#include "joint_matrix_hip_apply.hpp" +#include "joint_matrix_hip_copy.hpp" +#include "joint_matrix_hip_fill.hpp" +#include "joint_matrix_hip_mfma.hpp" + +// gfx942 (CDNA3) supports the same fp16/bf16/fp64 shapes as gfx90a, but its +// int8 MFMA uses larger K dimensions: 16x16x32 and 32x32x16. +template void matrix_mfma() { + hip_matrix_mfma(); + hip_matrix_mfma(); + hip_matrix_mfma(); + hip_matrix_mfma(); + hip_matrix_mfma(); + hip_matrix_mfma(); + hip_matrix_mfma(); + hip_matrix_mfma(); + hip_matrix_mfma(); + hip_matrix_mfma(); +} + +int main() { + matrix_mfma<1>(); + matrix_mfma<2>(); + matrix_mfma<3>(); + matrix_mfma<4>(); + + hip_matrix_copy(); + hip_matrix_copy(); + hip_matrix_copy(); + hip_matrix_copy(); + hip_matrix_copy(); + hip_matrix_copy(); + hip_matrix_copy(); + hip_matrix_copy(); + hip_matrix_copy(); + hip_matrix_copy(); + + hip_matrix_fill(); + hip_matrix_fill(); + hip_matrix_fill(); + hip_matrix_fill(); + hip_matrix_fill(); + + hip_matrix_apply(); + hip_matrix_apply(); + hip_matrix_apply(); + hip_matrix_apply(); + hip_matrix_apply(); +} diff --git a/sycl/test-e2e/Matrix/joint_matrix_hip_half_gfx942.cpp b/sycl/test-e2e/Matrix/joint_matrix_hip_half_gfx942.cpp new file mode 100644 index 0000000000000..a2ca215918884 --- /dev/null +++ b/sycl/test-e2e/Matrix/joint_matrix_hip_half_gfx942.cpp @@ -0,0 +1,44 @@ +//===---joint_matrix_hip_half_gfx942.cpp - DPC++ joint_matrix-------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// RUN: %clangxx -fsycl -fsycl-targets=amd_gpu_gfx942 %s -o %t.out +// RUN: %{run} %t.out + +// REQUIRES: target-amd +// REQUIRES: arch-amd_gpu_gfx942 +// REQUIRES: aspect-fp16 + +#include "joint_matrix_hip_apply.hpp" +#include "joint_matrix_hip_copy.hpp" +#include "joint_matrix_hip_fill.hpp" +#include "joint_matrix_hip_mfma.hpp" + +template void half_matrix_mfma() { + hip_matrix_mfma(); + hip_matrix_mfma(); + hip_matrix_mfma(); + hip_matrix_mfma(); +} + +int main() { + half_matrix_mfma<1>(); + half_matrix_mfma<2>(); + half_matrix_mfma<3>(); + half_matrix_mfma<4>(); + + hip_matrix_copy(); + hip_matrix_copy(); + hip_matrix_copy(); + hip_matrix_copy(); + + hip_matrix_fill(); + hip_matrix_fill(); + + hip_matrix_apply(); + hip_matrix_apply(); +} diff --git a/sycl/test-e2e/Matrix/runtime_query_hip_gfx942.cpp b/sycl/test-e2e/Matrix/runtime_query_hip_gfx942.cpp new file mode 100644 index 0000000000000..5f315eff5ef6b --- /dev/null +++ b/sycl/test-e2e/Matrix/runtime_query_hip_gfx942.cpp @@ -0,0 +1,68 @@ +//===---runtime_query_hip_gfx942.cpp - DPC++ joint_matrix-----------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// REQUIRES: target-amd +// REQUIRES: arch-amd_gpu_gfx942 +// RUN: %clangxx -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx942 %s -o %t.out +// RUN: %{run} %t.out + +#include +#include + +using namespace sycl::ext::oneapi::experimental::matrix; + +bool find_combination(const combination &comb, + const std::vector &expected_combinations) { + return std::find_if(expected_combinations.begin(), + expected_combinations.end(), + [&comb](const auto &expected_comb) { + return (comb.max_msize == expected_comb.max_msize && + comb.max_nsize == expected_comb.max_nsize && + comb.max_ksize == expected_comb.max_ksize && + comb.msize == expected_comb.msize && + comb.nsize == expected_comb.nsize && + comb.ksize == expected_comb.ksize && + comb.atype == expected_comb.atype && + comb.btype == expected_comb.btype && + comb.ctype == expected_comb.ctype && + comb.dtype == expected_comb.dtype); + }) != expected_combinations.end(); +} + +int main() { + std::vector expected_combinations = { + {0, 0, 0, 32, 32, 8, matrix_type::fp16, matrix_type::fp16, + matrix_type::fp32, matrix_type::fp32}, + {0, 0, 0, 16, 16, 16, matrix_type::fp16, matrix_type::fp16, + matrix_type::fp32, matrix_type::fp32}, + {0, 0, 0, 32, 32, 16, matrix_type::sint8, matrix_type::sint8, + matrix_type::sint32, matrix_type::sint32}, + {0, 0, 0, 16, 16, 32, matrix_type::sint8, matrix_type::sint8, + matrix_type::sint32, matrix_type::sint32}, + {0, 0, 0, 32, 32, 8, matrix_type::bf16, matrix_type::bf16, + matrix_type::fp32, matrix_type::fp32}, + {0, 0, 0, 16, 16, 16, matrix_type::bf16, matrix_type::bf16, + matrix_type::fp32, matrix_type::fp32}, + {0, 0, 0, 16, 16, 4, matrix_type::fp64, matrix_type::fp64, + matrix_type::fp64, matrix_type::fp64}}; + + sycl::queue q; + std::vector actual_combinations = + q.get_device() + .get_info(); + + assert(actual_combinations.size() == expected_combinations.size() && + "Number of combinations is not equal."); + + for (auto &comb : actual_combinations) { + assert(find_combination(comb, expected_combinations) && + "Some values in matrix runtime query for gfx942 are not expected."); + } + return 0; +} diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-bfloat16-float-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-bfloat16-float-test.cpp index a7409dd7426bb..1b8175c07471a 100644 --- a/sycl/test/check_device_code/hip/matrix/matrix-hip-bfloat16-float-test.cpp +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-bfloat16-float-test.cpp @@ -1,4 +1,5 @@ // REQUIRES: hip +// REQUIRES: hip-arch-gfx90a // RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s #include diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-double-double-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-double-double-test.cpp index f22ea888af2fe..8b4819252bb4e 100644 --- a/sycl/test/check_device_code/hip/matrix/matrix-hip-double-double-test.cpp +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-double-double-test.cpp @@ -1,4 +1,5 @@ // REQUIRES: hip +// REQUIRES: hip-arch-gfx90a // RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s #include diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-bfloat16-float-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-bfloat16-float-test.cpp new file mode 100644 index 0000000000000..30410c94bf500 --- /dev/null +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-bfloat16-float-test.cpp @@ -0,0 +1,65 @@ +// REQUIRES: hip +// REQUIRES: hip-arch-gfx942 +// RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s + +// gfx942 (CDNA3) uses the same bfloat16 MFMA instructions and shapes as gfx90a. + +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; +using sycl::ext::oneapi::bfloat16; + +SYCL_EXTERNAL [[sycl::reqd_work_group_size(1, 1, 64)]] void +row_row_m16n16k16(sycl::accessor + accA, + sycl::accessor + accB, + sycl::accessor + accC, + sycl::accessor + accD, + nd_item<2> item) { + sycl::sub_group sg = item.get_sub_group(); + + joint_matrix sub_c{}; + joint_matrix sub_a{}; + joint_matrix sub_b{}; + + // CHECK: tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> zeroinitializer, <4 x i16> zeroinitializer, <4 x float> zeroinitializer, i32 0, i32 0, i32 0) + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + joint_matrix_store(sg, sub_c, + accD.template get_multi_ptr(), 16, + layout::row_major); +} + +SYCL_EXTERNAL [[sycl::reqd_work_group_size(1, 1, 64)]] void +row_col_m32n32k8(sycl::accessor + accA, + sycl::accessor + accB, + sycl::accessor + accC, + sycl::accessor + accD, + nd_item<2> item) { + sycl::sub_group sg = item.get_sub_group(); + + joint_matrix sub_c{}; + joint_matrix sub_a{}; + joint_matrix sub_b{}; + + // CHECK: tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8bf16.1k(<4 x i16> zeroinitializer, <4 x i16> zeroinitializer, <16 x float> zeroinitializer, i32 0, i32 0, i32 0) + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + joint_matrix_store(sg, sub_c, + accD.template get_multi_ptr(), 32, + layout::row_major); +} diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-double-double-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-double-double-test.cpp new file mode 100644 index 0000000000000..b622f44e33c32 --- /dev/null +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-double-double-test.cpp @@ -0,0 +1,37 @@ +// REQUIRES: hip +// REQUIRES: hip-arch-gfx942 +// RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s + +// gfx942 (CDNA3) uses the same double MFMA instruction and shape as gfx90a. + +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; + +SYCL_EXTERNAL [[sycl::reqd_work_group_size(1, 1, 64)]] void +row_row_m16n16k4(sycl::accessor + accA, + sycl::accessor + accB, + sycl::accessor + accC, + sycl::accessor + accD, + nd_item<2> item) { + sycl::sub_group sg = item.get_sub_group(); + + joint_matrix sub_c{}; + joint_matrix sub_a{}; + joint_matrix sub_b{}; + + // CHECK: tail call <4 x double> @llvm.amdgcn.mfma.f64.16x16x4f64(double {{.*}}, double {{.*}}, <4 x double> zeroinitializer, i32 0, i32 0, i32 0) + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + joint_matrix_store(sg, sub_c, + accD.template get_multi_ptr(), 16, + layout::row_major); +} diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-half-float-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-half-float-test.cpp new file mode 100644 index 0000000000000..03be434db21ed --- /dev/null +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-half-float-test.cpp @@ -0,0 +1,64 @@ +// REQUIRES: hip +// REQUIRES: hip-arch-gfx942 +// RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s + +// gfx942 (CDNA3) uses the same half MFMA instructions and shapes as gfx90a. + +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; + +SYCL_EXTERNAL [[sycl::reqd_work_group_size(1, 1, 64)]] void +row_row_m16n16k16(sycl::accessor + accA, + sycl::accessor + accB, + sycl::accessor + accC, + sycl::accessor + accD, + nd_item<2> item) { + sycl::sub_group sg = item.get_sub_group(); + + joint_matrix sub_c{}; + joint_matrix sub_a{}; + joint_matrix sub_b{}; + + // CHECK: tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> zeroinitializer, i32 0, i32 0, i32 0) + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + joint_matrix_store(sg, sub_c, + accD.template get_multi_ptr(), 16, + layout::row_major); +} + +SYCL_EXTERNAL [[sycl::reqd_work_group_size(1, 1, 64)]] void +row_col_m32n32k8(sycl::accessor + accA, + sycl::accessor + accB, + sycl::accessor + accC, + sycl::accessor + accD, + nd_item<2> item) { + sycl::sub_group sg = item.get_sub_group(); + + joint_matrix sub_c{}; + joint_matrix sub_a{}; + joint_matrix sub_b{}; + + // CHECK: tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> zeroinitializer, i32 0, i32 0, i32 0) + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + joint_matrix_store(sg, sub_c, + accD.template get_multi_ptr(), 32, + layout::row_major); +} diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-int8-int32-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-int8-int32-test.cpp new file mode 100644 index 0000000000000..06d6a9211e172 --- /dev/null +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-int8-int32-test.cpp @@ -0,0 +1,65 @@ +// REQUIRES: hip +// REQUIRES: hip-arch-gfx942 +// RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s + +// gfx942 (CDNA3) int8 MFMA uses larger K dimensions than gfx90a: +// 16x16x32 and 32x32x16, with the A/B operands packed into i64 values. + +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; + +SYCL_EXTERNAL [[sycl::reqd_work_group_size(1, 1, 64)]] void +row_row_m16n16k32(sycl::accessor + accA, + sycl::accessor + accB, + sycl::accessor + accC, + sycl::accessor + accD, + nd_item<2> item) { + sycl::sub_group sg = item.get_sub_group(); + + joint_matrix sub_c{}; + joint_matrix sub_a{}; + joint_matrix sub_b{}; + + // CHECK: tail call <4 x i32> @llvm.amdgcn.mfma.i32.16x16x32.i8(i64 {{.*}}, i64 {{.*}}, <4 x i32> zeroinitializer, i32 0, i32 0, i32 0) + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + joint_matrix_store(sg, sub_c, + accD.template get_multi_ptr(), 16, + layout::row_major); +} + +SYCL_EXTERNAL [[sycl::reqd_work_group_size(1, 1, 64)]] void +row_col_m32n32k16(sycl::accessor + accA, + sycl::accessor + accB, + sycl::accessor + accC, + sycl::accessor + accD, + nd_item<2> item) { + sycl::sub_group sg = item.get_sub_group(); + + joint_matrix sub_c{}; + joint_matrix sub_a{}; + joint_matrix sub_b{}; + + // CHECK: tail call <16 x i32> @llvm.amdgcn.mfma.i32.32x32x16.i8(i64 {{.*}}, i64 {{.*}}, <16 x i32> zeroinitializer, i32 0, i32 0, i32 0) + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + joint_matrix_store(sg, sub_c, + accD.template get_multi_ptr(), 32, + layout::row_major); +} diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-half-float-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-half-float-test.cpp index e712f5c757892..ca030beff6d65 100644 --- a/sycl/test/check_device_code/hip/matrix/matrix-hip-half-float-test.cpp +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-half-float-test.cpp @@ -1,4 +1,5 @@ // REQUIRES: hip +// REQUIRES: hip-arch-gfx90a // RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s #include diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-int8-int32-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-int8-int32-test.cpp index 159bfff209131..db6e80004da23 100644 --- a/sycl/test/check_device_code/hip/matrix/matrix-hip-int8-int32-test.cpp +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-int8-int32-test.cpp @@ -1,4 +1,5 @@ // REQUIRES: hip +// REQUIRES: hip-arch-gfx90a // RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s #include diff --git a/sycl/test/lit.cfg.py b/sycl/test/lit.cfg.py index 599fe1f95e7a8..48f1c00fd2a8f 100644 --- a/sycl/test/lit.cfg.py +++ b/sycl/test/lit.cfg.py @@ -177,14 +177,30 @@ if "amdgcn-amd-amdhsa" in triple: llvm_config.with_system_environment("ROCM_PATH") config.available_features.add("hip") - # For AMD the specific GPU has to be specified with --offload-arch - if not any([f.startswith("--offload-arch") for f in additional_flags]): - # If the offload arch wasn't specified in SYCL_CLANG_EXTRA_FLAGS, - # hardcode it to gfx90a, this is fine because only compiler tests + # For AMD the specific GPU has to be specified with --offload-arch. The + # target arch used for every device compilation is configurable via the + # `amd_arch` lit param (default gfx90a), so the same suite can be run + # against new architectures without any harness changes, e.g.: + # llvm-lit --param SYCL_TRIPLE=amdgcn-amd-amdhsa --param amd_arch=gfx942 + # + # If SYCL_CLANG_EXTRA_FLAGS already provides --offload-arch, that value is + # honored instead. Either way, a `hip-arch-` feature is exposed so + # arch-specific tests (e.g. matrix codegen, whose emitted intrinsics differ + # per arch) can gate themselves with `REQUIRES: hip-arch-gfx942`. This means + # only the tests matching the selected arch run in a given invocation, which + # is the natural model since a single device compilation targets one arch. + offload_arch_flags = [ + f for f in additional_flags if f.startswith("--offload-arch=") + ] + if offload_arch_flags: + amd_arch = offload_arch_flags[-1].split("=", 1)[1] + else: + amd_arch = lit_config.params.get("amd_arch", "gfx90a") additional_flags += [ "-Xsycl-target-backend=amdgcn-amd-amdhsa", - "--offload-arch=gfx90a", + "--offload-arch=" + amd_arch, ] + config.available_features.add("hip-arch-" + amd_arch) config.sycl_headers_filter = lit_config.params.get("SYCL_HEADERS_FILTER", None) if config.sycl_headers_filter is not None: diff --git a/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp b/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp new file mode 100644 index 0000000000000..649bda8eea77c --- /dev/null +++ b/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp @@ -0,0 +1,28 @@ +// REQUIRES: hip +// RUN: %clangxx -fsycl -fsycl-targets=amd_gpu_gfx942 %s -o compile-query-hip + +#include +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental; +using namespace sycl::ext::oneapi::experimental::matrix; + +int main() { + // Compile-time query to validate the matrix parameters + using myparams = matrix_params; + + static_assert(myparams::M == 32); + static_assert(myparams::N == 32); + static_assert(myparams::K == 16); + + // Sizes-only compile-time query: types are given, generate default sizes + using myparams2 = matrix_params; + static_assert(myparams2::M == 16); + static_assert(myparams2::N == 16); + static_assert(myparams2::K == 32); + + return 0; +}; From 91d9d2805068ac106165f09cfe2d0b68cd04d490 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Mon, 20 Jul 2026 14:06:10 -0700 Subject: [PATCH 3/5] [SYCL][HIP] Extend joint_matrix compile-time support to gfx940/gfx941 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 --- .../ext/oneapi/matrix/static-query-use.hpp | 127 ++++++++++++++++++ .../matrix-hip-gfx942-bfloat16-float-test.cpp | 2 +- .../matrix-hip-gfx942-double-double-test.cpp | 2 +- .../matrix-hip-gfx942-half-float-test.cpp | 2 +- .../matrix-hip-gfx942-int8-int32-test.cpp | 2 +- .../matrix/hip/compile-query-hip-gfx942.cpp | 17 ++- 6 files changed, 143 insertions(+), 9 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp b/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp index 0c1ea1d2c536a..b2e15a5b1a3ff 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp @@ -664,6 +664,133 @@ struct matrix_params< using joint_matrix_d = joint_matrix; }; +////////////////////////////////////////////////////// +/// AMD Matrix Cores - GFX940 / GFX941 architectures /// +////////////////////////////////////////////////////// +// +// gfx940 and gfx941 are CDNA3 parts and share the exact same joint_matrix +// support as gfx942, so they reuse the gfx942 type/combination helpers. + +// Default-values query for gfx940: +template +struct matrix_params< + architecture::amd_gpu_gfx940, Ta, Tb, Tc, Td, 0, 0, 0, + typename std::enable_if_t<( + !std::is_same_v && !std::is_same_v && + !std::is_same_v && !std::is_same_v && + std::is_same_v && std::is_same_v)>> { + static_assert( + are_types_valid_amd_gfx942(), + "Invalid types for AMD gfx940, supported types are half, float, " + "int8_t, int32_t, double and bfloat16 "); + + // Default sizes for AMD gfx940 were chosen to represent a square matrix + static constexpr std::size_t M = 16; + static constexpr std::size_t N = 16; + static constexpr std::size_t K = std::is_same_v ? 4 + : std::is_same_v ? 32 + : 16; + + template + using joint_matrix_a = joint_matrix; + template + using joint_matrix_b = joint_matrix; + template + using joint_matrix_c = joint_matrix; + template + using joint_matrix_d = joint_matrix; +}; + +// Validation query for gfx940: +template +struct matrix_params< + architecture::amd_gpu_gfx940, Ta, Tb, Tc, Td, sM, sN, sK, + typename std::enable_if_t<( + !std::is_same_v && !std::is_same_v && + !std::is_same_v && !std::is_same_v && + std::is_same_v && std::is_same_v && sM != 0 && + sN != 0 && sK != 0)>> { + static_assert( + is_combination_valid_amd_gfx942(sM, sN, sK), + "Invalid parameters for AMD gfx940, query valid combinations " + "using: " + "q.get_device().get_info()"); + + static constexpr std::size_t M = sM; + static constexpr std::size_t N = sN; + static constexpr std::size_t K = sK; + + template + using joint_matrix_a = joint_matrix; + template + using joint_matrix_b = joint_matrix; + template + using joint_matrix_c = joint_matrix; + template + using joint_matrix_d = joint_matrix; +}; + +// Default-values query for gfx941: +template +struct matrix_params< + architecture::amd_gpu_gfx941, Ta, Tb, Tc, Td, 0, 0, 0, + typename std::enable_if_t<( + !std::is_same_v && !std::is_same_v && + !std::is_same_v && !std::is_same_v && + std::is_same_v && std::is_same_v)>> { + static_assert( + are_types_valid_amd_gfx942(), + "Invalid types for AMD gfx941, supported types are half, float, " + "int8_t, int32_t, double and bfloat16 "); + + // Default sizes for AMD gfx941 were chosen to represent a square matrix + static constexpr std::size_t M = 16; + static constexpr std::size_t N = 16; + static constexpr std::size_t K = std::is_same_v ? 4 + : std::is_same_v ? 32 + : 16; + + template + using joint_matrix_a = joint_matrix; + template + using joint_matrix_b = joint_matrix; + template + using joint_matrix_c = joint_matrix; + template + using joint_matrix_d = joint_matrix; +}; + +// Validation query for gfx941: +template +struct matrix_params< + architecture::amd_gpu_gfx941, Ta, Tb, Tc, Td, sM, sN, sK, + typename std::enable_if_t<( + !std::is_same_v && !std::is_same_v && + !std::is_same_v && !std::is_same_v && + std::is_same_v && std::is_same_v && sM != 0 && + sN != 0 && sK != 0)>> { + static_assert( + is_combination_valid_amd_gfx942(sM, sN, sK), + "Invalid parameters for AMD gfx941, query valid combinations " + "using: " + "q.get_device().get_info()"); + + static constexpr std::size_t M = sM; + static constexpr std::size_t N = sN; + static constexpr std::size_t K = sK; + + template + using joint_matrix_a = joint_matrix; + template + using joint_matrix_b = joint_matrix; + template + using joint_matrix_c = joint_matrix; + template + using joint_matrix_d = joint_matrix; +}; + ///////////////////////////////////////////////// /// CUDA Tensor Cores - sm70, sm72 and sm80 /// ///////////////////////////////////////////////// diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-bfloat16-float-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-bfloat16-float-test.cpp index 30410c94bf500..3692e10028dd9 100644 --- a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-bfloat16-float-test.cpp +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-bfloat16-float-test.cpp @@ -1,5 +1,5 @@ // REQUIRES: hip -// REQUIRES: hip-arch-gfx942 +// REQUIRES: hip-arch-gfx942 || hip-arch-gfx941 || hip-arch-gfx940 // RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s // gfx942 (CDNA3) uses the same bfloat16 MFMA instructions and shapes as gfx90a. diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-double-double-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-double-double-test.cpp index b622f44e33c32..59728f56c8927 100644 --- a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-double-double-test.cpp +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-double-double-test.cpp @@ -1,5 +1,5 @@ // REQUIRES: hip -// REQUIRES: hip-arch-gfx942 +// REQUIRES: hip-arch-gfx942 || hip-arch-gfx941 || hip-arch-gfx940 // RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s // gfx942 (CDNA3) uses the same double MFMA instruction and shape as gfx90a. diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-half-float-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-half-float-test.cpp index 03be434db21ed..148671149c4a9 100644 --- a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-half-float-test.cpp +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-half-float-test.cpp @@ -1,5 +1,5 @@ // REQUIRES: hip -// REQUIRES: hip-arch-gfx942 +// REQUIRES: hip-arch-gfx942 || hip-arch-gfx941 || hip-arch-gfx940 // RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s // gfx942 (CDNA3) uses the same half MFMA instructions and shapes as gfx90a. diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-int8-int32-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-int8-int32-test.cpp index 06d6a9211e172..1c6eb5de2930a 100644 --- a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-int8-int32-test.cpp +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-int8-int32-test.cpp @@ -1,5 +1,5 @@ // REQUIRES: hip -// REQUIRES: hip-arch-gfx942 +// REQUIRES: hip-arch-gfx942 || hip-arch-gfx941 || hip-arch-gfx940 // RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s // gfx942 (CDNA3) int8 MFMA uses larger K dimensions than gfx90a: diff --git a/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp b/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp index 649bda8eea77c..e1b424b8d2b04 100644 --- a/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp +++ b/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp @@ -8,21 +8,28 @@ using namespace sycl; using namespace sycl::ext::oneapi::experimental; using namespace sycl::ext::oneapi::experimental::matrix; -int main() { +// gfx940, gfx941 and gfx942 are CDNA3 parts and expose identical joint_matrix +// support, so the compile-time query must succeed for all three architectures. +template void check_arch() { // Compile-time query to validate the matrix parameters - using myparams = matrix_params; + using myparams = + matrix_params; static_assert(myparams::M == 32); static_assert(myparams::N == 32); static_assert(myparams::K == 16); // Sizes-only compile-time query: types are given, generate default sizes - using myparams2 = matrix_params; + using myparams2 = matrix_params; static_assert(myparams2::M == 16); static_assert(myparams2::N == 16); static_assert(myparams2::K == 32); +} + +int main() { + check_arch(); + check_arch(); + check_arch(); return 0; }; From 94a6347badf8b1af005c49eb750911c09ee91bab Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Tue, 21 Jul 2026 19:07:46 -0700 Subject: [PATCH 4/5] [SYCL][HIP] Add gfx942 joint_matrix float/float MFMA combinations 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 --- .../sycl/ext/oneapi/matrix/matrix-hip.hpp | 21 +++++- .../ext/oneapi/matrix/static-query-use.hpp | 28 +++++--- .../matrix-hip-gfx942-float-float-test.cpp | 64 +++++++++++++++++++ .../matrix/hip/compile-query-hip-gfx942.cpp | 9 +++ 4 files changed, 112 insertions(+), 10 deletions(-) create mode 100644 sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-float-float-test.cpp diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp index e66ec2ba73651..21100368c8a14 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp @@ -83,6 +83,11 @@ __SYCL_JOINT_MATRIX_OVERLOAD_ARR(half, b, 16, 16, 4) __SYCL_JOINT_MATRIX_OVERLOAD_ARR(half, a, 32, 8, 4) __SYCL_JOINT_MATRIX_OVERLOAD_ARR(half, b, 8, 32, 4) +__SYCL_JOINT_MATRIX_OVERLOAD_ARR(float, a, 16, 4, 1) +__SYCL_JOINT_MATRIX_OVERLOAD_ARR(float, b, 4, 16, 1) +__SYCL_JOINT_MATRIX_OVERLOAD_ARR(float, a, 32, 2, 1) +__SYCL_JOINT_MATRIX_OVERLOAD_ARR(float, b, 2, 32, 1) + __SYCL_JOINT_MATRIX_OVERLOAD_ARR(double, a, 16, 4, 1) __SYCL_JOINT_MATRIX_OVERLOAD_ARR(double, b, 4, 16, 1) @@ -379,7 +384,21 @@ void joint_matrix_mad_hip( sycl::ext::oneapi::experimental::matrix::layout::dynamic> &C) { #if defined(__gfx90a__) || defined(__gfx940__) || defined(__gfx941__) || \ defined(__gfx942__) - if constexpr (std::is_same_v) { + if constexpr (std::is_same_v) { +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + if constexpr (M == 16 && N == 16) { + auto result = __builtin_amdgcn_mfma_f32_16x16x4f32( + A.wi_marray[0], B.wi_marray[0], + *reinterpret_cast(&C.wi_marray), 0, 0, 0); + std::memcpy(&D.wi_marray, &result, 4 * sizeof(float)); + } else if constexpr (M == 32 && N == 32) { + auto result = __builtin_amdgcn_mfma_f32_32x32x2f32( + A.wi_marray[0], B.wi_marray[0], + *reinterpret_cast(&C.wi_marray), 0, 0, 0); + std::memcpy(&D.wi_marray, &result, 16 * sizeof(float)); + } +#endif + } else if constexpr (std::is_same_v) { if constexpr (M == 16 && N == 16) { auto result = __builtin_amdgcn_mfma_f32_16x16x16f16( *reinterpret_cast(&A.wi_marray), diff --git a/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp b/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp index b2e15a5b1a3ff..b2af4efe4b775 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp @@ -584,6 +584,9 @@ constexpr bool is_combination_valid_amd_gfx942(size_t sM, size_t sN, return (std::is_same_v && std::is_same_v && ((sM == 32 && sN == 32 && sK == 8) || (sM == 16 && sN == 16 && sK == 16))) || + (std::is_same_v && std::is_same_v && + ((sM == 32 && sN == 32 && sK == 2) || + (sM == 16 && sN == 16 && sK == 4))) || (std::is_same_v && std::is_same_v && ((sM == 32 && sN == 32 && sK == 16) || (sM == 16 && sN == 16 && sK == 32))) || @@ -597,6 +600,7 @@ constexpr bool is_combination_valid_amd_gfx942(size_t sM, size_t sN, template constexpr bool are_types_valid_amd_gfx942() { return (std::is_same_v && std::is_same_v) || + (std::is_same_v && std::is_same_v) || (std::is_same_v && std::is_same_v) || (std::is_same_v && std::is_same_v) || (std::is_same_v && std::is_same_v); @@ -619,9 +623,11 @@ struct matrix_params< // Default sizes for AMD gfx942 were chosen to represent a square matrix static constexpr std::size_t M = 16; static constexpr std::size_t N = 16; - static constexpr std::size_t K = std::is_same_v ? 4 - : std::is_same_v ? 32 - : 16; + static constexpr std::size_t K = + std::is_same_v ? 32 + : (std::is_same_v || std::is_same_v) + ? 16 + : 4; template using joint_matrix_a = joint_matrix; @@ -687,9 +693,11 @@ struct matrix_params< // Default sizes for AMD gfx940 were chosen to represent a square matrix static constexpr std::size_t M = 16; static constexpr std::size_t N = 16; - static constexpr std::size_t K = std::is_same_v ? 4 - : std::is_same_v ? 32 - : 16; + static constexpr std::size_t K = + std::is_same_v ? 32 + : (std::is_same_v || std::is_same_v) + ? 16 + : 4; template using joint_matrix_a = joint_matrix; @@ -747,9 +755,11 @@ struct matrix_params< // Default sizes for AMD gfx941 were chosen to represent a square matrix static constexpr std::size_t M = 16; static constexpr std::size_t N = 16; - static constexpr std::size_t K = std::is_same_v ? 4 - : std::is_same_v ? 32 - : 16; + static constexpr std::size_t K = + std::is_same_v ? 32 + : (std::is_same_v || std::is_same_v) + ? 16 + : 4; template using joint_matrix_a = joint_matrix; diff --git a/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-float-float-test.cpp b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-float-float-test.cpp new file mode 100644 index 0000000000000..835605342798e --- /dev/null +++ b/sycl/test/check_device_code/hip/matrix/matrix-hip-gfx942-float-float-test.cpp @@ -0,0 +1,64 @@ +// REQUIRES: hip +// REQUIRES: hip-arch-gfx942 || hip-arch-gfx941 || hip-arch-gfx940 +// RUN: %clangxx -fsycl-device-only -fsycl-targets=amdgcn-amd-amdhsa -S %s -o -| FileCheck %s + +// CDNA3 supports one-block F32 MFMA shapes: 16x16x4 and 32x32x2. + +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental::matrix; + +SYCL_EXTERNAL [[sycl::reqd_work_group_size(1, 1, 64)]] void +row_row_m16n16k4(sycl::accessor + accA, + sycl::accessor + accB, + sycl::accessor + accC, + sycl::accessor + accD, + nd_item<2> item) { + sycl::sub_group sg = item.get_sub_group(); + + joint_matrix sub_c{}; + joint_matrix sub_a{}; + joint_matrix sub_b{}; + + // CHECK: tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x4f32(float {{.*}}, float {{.*}}, <4 x float> zeroinitializer, i32 0, i32 0, i32 0) + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + joint_matrix_store(sg, sub_c, + accD.template get_multi_ptr(), 16, + layout::row_major); +} + +SYCL_EXTERNAL [[sycl::reqd_work_group_size(1, 1, 64)]] void +row_col_m32n32k2(sycl::accessor + accA, + sycl::accessor + accB, + sycl::accessor + accC, + sycl::accessor + accD, + nd_item<2> item) { + sycl::sub_group sg = item.get_sub_group(); + + joint_matrix sub_c{}; + joint_matrix sub_a{}; + joint_matrix sub_b{}; + + // CHECK: tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x2f32(float {{.*}}, float {{.*}}, <16 x float> zeroinitializer, i32 0, i32 0, i32 0) + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + joint_matrix_store(sg, sub_c, + accD.template get_multi_ptr(), 32, + layout::row_major); +} diff --git a/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp b/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp index e1b424b8d2b04..6f0cf6120c00e 100644 --- a/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp +++ b/sycl/test/matrix/hip/compile-query-hip-gfx942.cpp @@ -14,16 +14,25 @@ template void check_arch() { // Compile-time query to validate the matrix parameters using myparams = matrix_params; + using floatparams = + matrix_params; static_assert(myparams::M == 32); static_assert(myparams::N == 32); static_assert(myparams::K == 16); + static_assert(floatparams::M == 32); + static_assert(floatparams::N == 32); + static_assert(floatparams::K == 2); // Sizes-only compile-time query: types are given, generate default sizes using myparams2 = matrix_params; + using floatparams2 = matrix_params; static_assert(myparams2::M == 16); static_assert(myparams2::N == 16); static_assert(myparams2::K == 32); + static_assert(floatparams2::M == 16); + static_assert(floatparams2::N == 16); + static_assert(floatparams2::K == 4); } int main() { From 290d359ada70fb2715c05b320d38cf6337ff6601 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Wed, 22 Jul 2026 09:26:01 -0700 Subject: [PATCH 5/5] [SYCL][HIP] Fix clang-format issues Co-authored-by: Cursor --- clang/lib/CodeGen/CodeGenModule.cpp | 11 ++++++++++ .../sycl/ext/oneapi/matrix/matrix-hip.hpp | 8 +++---- .../ext/oneapi/matrix/static-query-use.hpp | 21 ++++++++----------- 3 files changed, 24 insertions(+), 16 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 036a0de7be2b8..71033821d36ee 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1255,6 +1255,17 @@ void CodeGenModule::Release() { : "buffered"); getModule().addModuleFlag(llvm::Module::Error, "amdgpu_printf_kind", MDStr); + } else if (LangOpts.SYCLIsDevice) { + // SYCL device code on AMDGPU always uses the buffered printf scheme: the + // hostcall scheme relies on ROCm's hostcall runtime service (and PCIe + // atomics), which the SYCL/UR HIP runtime does not set up. The buffered + // scheme only needs the printf buffer implicit kernarg, which the HIP + // runtime allocates and parses automatically from the code object's + // printf metadata. This flag tells AMDGPUPrintfRuntimeBinding to lower + // printf using the buffered layout expected by the runtime. + getModule().addModuleFlag( + llvm::Module::Error, "amdgpu_printf_kind", + llvm::MDString::get(getLLVMContext(), "buffered")); } } diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp index 21100368c8a14..67dfc8e319b51 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-hip.hpp @@ -236,10 +236,10 @@ void load_multiplicand_hip(joint_matrix_hip &res, // dimension (rows M for the A multiplicand, columns N for the B multiplicand) // and the shared contraction dimension K. Whether the free dimension is the // strided one in memory depends on BOTH the use and the layout: - // - A row_major: element (m,k) at m*stride + k -> free dim (m) is strided. - // - A col_major: element (m,k) at k*stride + m -> free dim (m) is contiguous. - // - B row_major: element (k,n) at k*stride + n -> free dim (n) is contiguous. - // - B col_major: element (k,n) at n*stride + k -> free dim (n) is strided. + // - A row_major: (m,k) at m*stride + k -> free dim (m) is strided. + // - A col_major: (m,k) at k*stride + m -> free dim (m) is contiguous. + // - B row_major: (k,n) at k*stride + n -> free dim (n) is contiguous. + // - B col_major: (k,n) at n*stride + k -> free dim (n) is strided. // So the strided-free-dim access pattern applies to (A, row_major) and // (B, col_major); the contiguous-free-dim pattern applies to (A, col_major) // and (B, row_major). Selecting the pattern from this combined condition diff --git a/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp b/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp index b2af4efe4b775..4d3f452323fea 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/static-query-use.hpp @@ -624,10 +624,9 @@ struct matrix_params< static constexpr std::size_t M = 16; static constexpr std::size_t N = 16; static constexpr std::size_t K = - std::is_same_v ? 32 - : (std::is_same_v || std::is_same_v) - ? 16 - : 4; + std::is_same_v ? 32 + : (std::is_same_v || std::is_same_v) ? 16 + : 4; template using joint_matrix_a = joint_matrix; @@ -694,10 +693,9 @@ struct matrix_params< static constexpr std::size_t M = 16; static constexpr std::size_t N = 16; static constexpr std::size_t K = - std::is_same_v ? 32 - : (std::is_same_v || std::is_same_v) - ? 16 - : 4; + std::is_same_v ? 32 + : (std::is_same_v || std::is_same_v) ? 16 + : 4; template using joint_matrix_a = joint_matrix; @@ -756,10 +754,9 @@ struct matrix_params< static constexpr std::size_t M = 16; static constexpr std::size_t N = 16; static constexpr std::size_t K = - std::is_same_v ? 32 - : (std::is_same_v || std::is_same_v) - ? 16 - : 4; + std::is_same_v ? 32 + : (std::is_same_v || std::is_same_v) ? 16 + : 4; template using joint_matrix_a = joint_matrix;