From 43951863151221401881d2d5c2a75d20b2b44da5 Mon Sep 17 00:00:00 2001 From: Dounia Khaldi Date: Mon, 31 Aug 2026 16:27:29 +0000 Subject: [PATCH 1/6] [SYCL][JM] Add support for FP8 and FP4 for CRI --- clang/lib/CodeGen/CodeGenTypes.cpp | 7 + .../include/sycl/__spirv/spirv_ops_matrix.hpp | 63 +++++ sycl/include/sycl/__spirv/spirv_types.hpp | 8 +- .../oneapi/matrix/matrix-unified-utils.hpp | 56 ++++- .../sycl/ext/oneapi/matrix/matrix-unified.hpp | 71 ++++++ .../sycl/ext/oneapi/matrix/query-types.hpp | 24 +- sycl/include/sycl/sycl.hpp | 2 + sycl/source/detail/device_impl.hpp | 26 +++ .../program_manager/program_manager.cpp | 9 + sycl/test-e2e/Matrix/Inputs/common.hpp | 156 +++++++++++++ .../Inputs/element_wise_all_ops_impl.hpp | 179 ++++++++++++++ .../Inputs/joint_matrix_float4_impl.hpp | 221 ++++++++++++++++++ .../Inputs/joint_matrix_float8_impl.hpp | 185 +++++++++++++++ .../Matrix/SG32/element_wise_all_ops.cpp | 2 +- .../Matrix/SG32/joint_matrix_float4.cpp | 19 ++ .../Matrix/SG32/joint_matrix_float8.cpp | 15 ++ sycl/test-e2e/Matrix/element_wise_all_ops.cpp | 2 +- sycl/test-e2e/Matrix/joint_matrix_float4.cpp | 18 ++ sycl/test-e2e/Matrix/joint_matrix_float8.cpp | 14 ++ 19 files changed, 1070 insertions(+), 7 deletions(-) create mode 100644 sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp create mode 100644 sycl/test-e2e/Matrix/Inputs/joint_matrix_float8_impl.hpp create mode 100644 sycl/test-e2e/Matrix/SG32/joint_matrix_float4.cpp create mode 100644 sycl/test-e2e/Matrix/SG32/joint_matrix_float8.cpp create mode 100644 sycl/test-e2e/Matrix/joint_matrix_float4.cpp create mode 100644 sycl/test-e2e/Matrix/joint_matrix_float8.cpp diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index fe005cb1b0efb..da81390598ea7 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -399,6 +399,13 @@ llvm::Type *CodeGenTypes::ConvertSPVCooperativeMatrixType(RecordDecl *RD) { CompTy = llvm::Type::getFloatTy(getLLVMContext()); } else if (LlvmTyName == "bfloat16") { CompTy = llvm::Type::getInt16Ty(getLLVMContext()); + // The 8-bit float types are class templates parameterized by the packing + // width (fp8_e5m2_x), so match on the class name without the width. + } else if (LlvmTyName == "fp8_e5m2_x" || LlvmTyName == "fp8_e4m3_x" || + LlvmTyName == "fp8_e8m0_x") { + CompTy = llvm::Type::getInt8Ty(getLLVMContext()); + } else if (LlvmTyName == "fp4_e2m1_x") { + CompTy = llvm::Type::getIntNTy(getLLVMContext(), 4); } else { llvm_unreachable("Wrong matrix base type!"); } diff --git a/sycl/include/sycl/__spirv/spirv_ops_matrix.hpp b/sycl/include/sycl/__spirv/spirv_ops_matrix.hpp index 0bd42ffca3081..8086ebb798f96 100644 --- a/sycl/include/sycl/__spirv/spirv_ops_matrix.hpp +++ b/sycl/include/sycl/__spirv/spirv_ops_matrix.hpp @@ -151,4 +151,67 @@ extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixPrefetchINTEL( T *Ptr, uint32_t NumRows, uint32_t NumCols, unsigned int CacheLevel, __spv::MatrixLayout Layout, size_t Stride); +// FP4E2M1 Upconversion +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_ConvertFP4E2M1ToHF16INTEL( + __spv::__spirv_CooperativeMatrixKHR *Object); +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_ConvertFP4E2M1ToBF16INTEL( + __spv::__spirv_CooperativeMatrixKHR *Object); +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_ConvertFP4E2M1ToHF8INTEL( + __spv::__spirv_CooperativeMatrixKHR *Object); +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_ConvertFP4E2M1ToBF8INTEL( + __spv::__spirv_CooperativeMatrixKHR *Object); +// FP4E2M1 down conversion +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_ConvertHF16ToFP4E2M1INTEL( + __spv::__spirv_CooperativeMatrixKHR *Object); +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_ConvertBF16ToFP4E2M1INTEL( + __spv::__spirv_CooperativeMatrixKHR *Object); + +template +extern __DPCPP_SYCL_EXTERNAL + __spv::__spirv_CooperativeMatrixKHR * + __spirv_CooperativeMatrixMulAddScaledINTEL( + __spv::__spirv_CooperativeMatrixKHR *A, + __spv::__spirv_CooperativeMatrixKHR *B, + __spv::__spirv_CooperativeMatrixKHR *C, + __spv::__spirv_CooperativeMatrixKHR *Ascale, + __spv::__spirv_CooperativeMatrixKHR *Bscale, + size_t Operands = 0); + #endif diff --git a/sycl/include/sycl/__spirv/spirv_types.hpp b/sycl/include/sycl/__spirv/spirv_types.hpp index 23a905a0b293d..31c93c371e522 100644 --- a/sycl/include/sycl/__spirv/spirv_types.hpp +++ b/sycl/include/sycl/__spirv/spirv_types.hpp @@ -130,7 +130,13 @@ enum class MatrixOperands : uint32_t { MatrixAAndBTF32ComponentsINTEL = 0x20, MatrixAAndBBFloat16ComponentsINTEL = 0x40, MatrixCBFloat16ComponentsINTEL = 0x80, - MatrixResultBFloat16ComponentsINTEL = 0x100 + MatrixResultBFloat16ComponentsINTEL = 0x100, + MatrixABFloat8ComponentsINTEL = 0x400, + MatrixBBFloat8ComponentsINTEL = 0x800, + MatrixAHFloat8ComponentsINTEL = 0x1000, + MatrixBHFloat8ComponentsINTEL = 0x2000, + MatrixAFP4S1E2M1ComponentsINTEL = 0x10000, + MatrixBFP4S1E2M1ComponentsINTEL = 0x20000 }; template // std::string_view #include // __spv namespace #include // bfloat16 -#include // std::pair +#include // fp4_e2m1_x +#include // fp8_e5m2, fp8_e4m3 +#include // std::pair namespace sycl { inline namespace _V1 { @@ -21,7 +23,7 @@ namespace oneapi { namespace experimental { namespace matrix { -enum class use { a, b, accumulator }; +enum class use { a, b, accumulator, scale }; enum class layout { row_major = 0, @@ -49,6 +51,7 @@ constexpr UseToUseStringPair UseToUseStringMap[] = { {ext::oneapi::experimental::matrix::use::a, "use::a"}, {ext::oneapi::experimental::matrix::use::b, "use::b"}, {ext::oneapi::experimental::matrix::use::accumulator, "use::accumulator"}, + {ext::oneapi::experimental::matrix::use::scale, "use::scale"}, }; constexpr const char * @@ -85,6 +88,13 @@ extern "C" constexpr __spv::MatrixLayout joint_matrix_layout_to_spv( } } +// fp4_e2m1_x is a packed type: any supported packing width N designates the +// same E2M1 element format, so matrix queries match on the format, not on N. +template struct is_fp4_e2m1 : std::false_type {}; +template +struct is_fp4_e2m1> + : std::true_type {}; + template constexpr uint32_t CalculateMatrixOperand() { uint32_t returnValue = 0x00; @@ -104,6 +114,48 @@ constexpr uint32_t CalculateMatrixOperand() { if constexpr (std::is_signed::value) returnValue += static_cast( __spv::MatrixOperands::MatrixBSignedComponentsKHR); + if constexpr (std::is_same< + Ta, sycl::ext::oneapi::experimental::fp8_e5m2>::value && + std::is_same< + Tb, sycl::ext::oneapi::experimental::fp8_e5m2>::value && + std::is_same::value) + returnValue += static_cast( + __spv::MatrixOperands::MatrixABFloat8ComponentsINTEL) + + static_cast( + __spv::MatrixOperands::MatrixBBFloat8ComponentsINTEL); + if constexpr (std::is_same< + Ta, sycl::ext::oneapi::experimental::fp8_e5m2>::value && + std::is_same< + Tb, sycl::ext::oneapi::experimental::fp8_e4m3>::value && + std::is_same::value) + returnValue += static_cast( + __spv::MatrixOperands::MatrixABFloat8ComponentsINTEL) + + static_cast( + __spv::MatrixOperands::MatrixBHFloat8ComponentsINTEL); + if constexpr (std::is_same< + Ta, sycl::ext::oneapi::experimental::fp8_e4m3>::value && + std::is_same< + Tb, sycl::ext::oneapi::experimental::fp8_e5m2>::value && + std::is_same::value) + returnValue += static_cast( + __spv::MatrixOperands::MatrixAHFloat8ComponentsINTEL) + + static_cast( + __spv::MatrixOperands::MatrixBBFloat8ComponentsINTEL); + if constexpr (std::is_same< + Ta, sycl::ext::oneapi::experimental::fp8_e4m3>::value && + std::is_same< + Tb, sycl::ext::oneapi::experimental::fp8_e4m3>::value && + std::is_same::value) + returnValue += static_cast( + __spv::MatrixOperands::MatrixAHFloat8ComponentsINTEL) + + static_cast( + __spv::MatrixOperands::MatrixBHFloat8ComponentsINTEL); + if constexpr (sycl::detail::is_fp4_e2m1::value && + sycl::detail::is_fp4_e2m1::value) + returnValue += static_cast( + __spv::MatrixOperands::MatrixAFP4S1E2M1ComponentsINTEL) + + static_cast( + __spv::MatrixOperands::MatrixBFP4S1E2M1ComponentsINTEL); return returnValue; } diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp index f107aabf39115..22642e821719e 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp @@ -576,6 +576,77 @@ joint_matrix_prefetch(Group sg, T *Ptr, size_t stride, #endif // defined(__SYCL_DEVICE_ONLY__) } +template +inline __SYCL_ALWAYS_INLINE void +joint_matrix_convert(Group, + const joint_matrix &src, + joint_matrix &dst) { +#if defined(__SYCL_DEVICE_ONLY__) + // FP4E2M1 Upconversion + if constexpr (sycl::detail::is_fp4_e2m1::value) { + if constexpr (std::is_same::value) + dst.spvm = __spirv_ConvertFP4E2M1ToHF16INTEL(src.spvm); + else if constexpr (std::is_same::value) + dst.spvm = __spirv_ConvertFP4E2M1ToBF16INTEL(src.spvm); + else if constexpr (std::is_same< + To, + sycl::ext::oneapi::experimental::fp8_e4m3>::value) + dst.spvm = __spirv_ConvertFP4E2M1ToHF8INTEL(src.spvm); + else if constexpr (std::is_same< + To, + sycl::ext::oneapi::experimental::fp8_e5m2>::value) + dst.spvm = __spirv_ConvertFP4E2M1ToBF8INTEL(src.spvm); + } + // FP4E2M1 down conversion + else if constexpr (sycl::detail::is_fp4_e2m1::value) { + if constexpr (std::is_same::value) + dst.spvm = __spirv_ConvertHF16ToFP4E2M1INTEL(src.spvm); + else if constexpr (std::is_same::value) + dst.spvm = __spirv_ConvertBF16ToFP4E2M1INTEL(src.spvm); + } +#else + std::ignore = src; + std::ignore = dst; + std::ignore = Layout; + throw exception(make_error_code(errc::runtime), + "joint_matrix_convert is not supported on host."); +#endif // defined(__SYCL_DEVICE_ONLY__) +} + +template +inline __SYCL_ALWAYS_INLINE void joint_matrix_bmad( + Group, + joint_matrix &D, + const joint_matrix &A, + const joint_matrix &B, + const joint_matrix &Ascale, + const joint_matrix &Bscale, + const joint_matrix + &C) { +#if defined(__SYCL_DEVICE_ONLY__) + constexpr uint32_t MatrixOperand = + sycl::detail::CalculateMatrixOperand(); + D.spvm = __spirv_CooperativeMatrixMulAddScaledINTEL( + A.spvm, B.spvm, C.spvm, Ascale.spvm, Bscale.spvm, MatrixOperand); +#else + std::ignore = A; + std::ignore = B; + std::ignore = Ascale; + std::ignore = Bscale; + std::ignore = C; + std::ignore = D; + throw exception(make_error_code(errc::runtime), + "joint matrix is not supported on host."); +#endif // defined(__SYCL_DEVICE_ONLY__) +} + } // namespace matrix } // namespace experimental } // namespace oneapi diff --git a/sycl/include/sycl/ext/oneapi/matrix/query-types.hpp b/sycl/include/sycl/ext/oneapi/matrix/query-types.hpp index b117a07b350b9..d366548dd3c49 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/query-types.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/query-types.hpp @@ -31,7 +31,10 @@ enum class matrix_type { uint8, uint16, uint32, - uint64 + uint64, + fp8_e5m2, + fp8_e4m3, + fp4_e2m1 }; struct combination { @@ -64,7 +67,12 @@ struct matrix_combinations // Type to matrix type string conversion used in compile-time namespace detail { template constexpr const char *convertTypeToMatrixTypeString() { - return ""; + // fp4_e2m1_x is a class template, so it cannot be handled by an explicit + // function template specialization; match on the element format instead. + if constexpr (is_fp4_e2m1::value) + return "matrix_type::fp4_e2m1"; + else + return ""; } template <> constexpr const char * @@ -79,6 +87,18 @@ constexpr const char *convertTypeToMatrixTypeString< sycl::ext::oneapi::experimental::matrix::precision::tf32>() { return "matrix_type::tf32"; } +template <> +constexpr const char * +convertTypeToMatrixTypeString() { + return "matrix_type::fp8_e5m2"; +} + +template <> +constexpr const char * +convertTypeToMatrixTypeString() { + return "matrix_type::fp8_e4m3"; +} + template <> constexpr const char *convertTypeToMatrixTypeString() { return "matrix_type::fp32"; } diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index d133824184470..a4cb9f99481e4 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -150,6 +150,8 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.") #include #include #include +#include +#include #include #include #include diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index ea45d51058753..99a61ee1a3c0c 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -2206,6 +2206,32 @@ class device_impl { {8, 0, 0, 0, 16, 8, matrix_type::tf32, matrix_type::tf32, matrix_type::fp32, matrix_type::fp32}, }; + // fp8 and fp4_e2m1 are supported starting with Crescent Island. + if (architecture::intel_gpu_cri == DeviceArch) + pvc_combs.insert( + pvc_combs.end(), + { + {8, 0, 0, 0, 16, 32, matrix_type::fp8_e5m2, + matrix_type::fp8_e5m2, matrix_type::bf16, matrix_type::bf16}, + {8, 0, 0, 0, 16, 32, matrix_type::fp8_e5m2, + matrix_type::fp8_e4m3, matrix_type::bf16, matrix_type::bf16}, + {8, 0, 0, 0, 16, 32, matrix_type::fp8_e4m3, + matrix_type::fp8_e5m2, matrix_type::bf16, matrix_type::bf16}, + {8, 0, 0, 0, 16, 32, matrix_type::fp8_e4m3, + matrix_type::fp8_e4m3, matrix_type::bf16, matrix_type::bf16}, + {8, 0, 0, 0, 16, 32, matrix_type::fp8_e5m2, + matrix_type::fp8_e5m2, matrix_type::fp32, matrix_type::fp32}, + {8, 0, 0, 0, 16, 32, matrix_type::fp8_e5m2, + matrix_type::fp8_e4m3, matrix_type::fp32, matrix_type::fp32}, + {8, 0, 0, 0, 16, 32, matrix_type::fp8_e4m3, + matrix_type::fp8_e5m2, matrix_type::fp32, matrix_type::fp32}, + {8, 0, 0, 0, 16, 32, matrix_type::fp8_e4m3, + matrix_type::fp8_e4m3, matrix_type::fp32, matrix_type::fp32}, + {0, 0, 0, 8, 16, 32, matrix_type::fp4_e2m1, + matrix_type::fp4_e2m1, matrix_type::fp32, matrix_type::fp32}, + {0, 0, 0, 8, 16, 32, matrix_type::fp4_e2m1, + matrix_type::fp4_e2m1, matrix_type::bf16, matrix_type::bf16}, + }); return pvc_combs; } else if ((architecture::intel_gpu_dg2_g10 == DeviceArch) || (architecture::intel_gpu_dg2_g11 == DeviceArch) || diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 5ca89927f5078..b73d1e349efa1 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -3105,6 +3105,12 @@ convertMatrixTypeStringMatrixTypeEnumValue( return matrix_ext::matrix_type::uint32; else if ("uint64" == MatrixTypeStringView) return matrix_ext::matrix_type::uint64; + else if ("fp8_e5m2" == MatrixTypeStringView) + return matrix_ext::matrix_type::fp8_e5m2; + else if ("fp8_e4m3" == MatrixTypeStringView) + return matrix_ext::matrix_type::fp8_e4m3; + else if ("fp4_e2m1" == MatrixTypeStringView) + return matrix_ext::matrix_type::fp4_e2m1; return std::nullopt; } @@ -3188,6 +3194,9 @@ std::optional checkDevSupportJointMatrix( Combination.nsize); break; } + case matrix_ext::use::scale: + // TODO as part of a new query + break; } // early exit if we have a match diff --git a/sycl/test-e2e/Matrix/Inputs/common.hpp b/sycl/test-e2e/Matrix/Inputs/common.hpp index f87cbfb992505..dbe42a5f85f4c 100644 --- a/sycl/test-e2e/Matrix/Inputs/common.hpp +++ b/sycl/test-e2e/Matrix/Inputs/common.hpp @@ -177,6 +177,162 @@ void matrix_copy(unsigned int rows, unsigned int cols, T *src, T *dst) { } } +// 8-bit float types: they pack a single element and convert to/from half. +template +constexpr bool is_fp8_type_v = std::is_same_v || + std::is_same_v || + std::is_same_v; + +// 4-bit types: they pack numElems elements and convert to/from +// marray. +template +constexpr bool is_4bit_type_v = std::is_same_v>; + +// numElems is the packing factor for the 4bits types. It can be 2 or 8 +// src is const so that Ts is deduced without a cv-qualifier when it comes from +// a read-only accessor: the is_*_type_v checks below match unqualified types. +template +void matrix_copy(queue q, unsigned int rows, unsigned int cols, const Ts *src, + Td *dst) { + q.single_task([=]() { + for (unsigned int i = 0; i < rows; i++) { + for (unsigned int j = 0; j < cols; j++) { + if constexpr (std::is_same_v && is_fp8_type_v) + dst[i * cols + j] = (Td)src[i * cols + j]; + else if constexpr (std::is_same_v && + is_fp8_type_v) + dst[i * cols + j] = src[i * cols + j]; + else if constexpr (std::is_same_v && + is_4bit_type_v) { + // src[i][j] packs numElems consecutive values, so it expands to + // dst[i][j * numElems .. j * numElems + numElems - 1]. + marray mval = (marray)src[i * cols + j]; + for (unsigned int p = 0; p < numElems; p++) + dst[i * cols * numElems + j * numElems + p] = mval[p]; + } else if constexpr (std::is_same_v && + is_4bit_type_v) { + marray mval; + for (unsigned int p = 0; p < numElems; p++) + mval[p] = src[i * cols * numElems + j * numElems + p]; + // The open 4-bit types only provide explicit conversions from + // marray, so construct rather than assign. + dst[i * cols + j] = Td(mval); + } else + assert(false && "Unsupported type in matrix_copy."); + } + } + }).wait(); +} + +template +Tin scalar_truncate_fraction_bits(Tin number) { + static_assert(std::is_same_v && + "Unsupported input type in scalar_truncate_fraction_bits"); + + union { + Tin as_input; + uint16_t as_integer; + } caster = {number}; + + if constexpr (std::is_same_v) { + caster.as_integer &= 0xFF00; + } else if constexpr (std::is_same_v) { + caster.as_integer &= 0xFF80; + } else + static_assert( + false && + "Unsupported precision type in scalar_truncate_fraction_bits."); + + return caster.as_input; +} + +template +void matrix_truncate_fraction_bits(unsigned int rows, unsigned int cols, + Tin *mat) { + for (unsigned int i = 0; i < rows; i++) { + for (unsigned int j = 0; j < cols; j++) { + mat[i * cols + j] = + scalar_truncate_fraction_bits(mat[i * cols + j]); + } + } +} + +template +void matrix_fill(queue q, unsigned int rows, unsigned int cols, T *src, F op) { + q.single_task([=]() { + for (unsigned int i = 0; i < rows; i++) { + for (unsigned int j = 0; j < cols; j++) { + if constexpr (is_fp8_type_v) + src[i * cols + j] = op(i, j); + else + assert(false && "Unsupported type in matrix_fill on device."); + } + } + }).wait(); +} + +// Scaling block size is what is used to calculate the scales. Its values is +// always 32 +constexpr size_t ScalingBlock = 32; + +template +void scaled_matrix_multiply_ref(Ta *A, Tb *B, syclex::fp8_e8m0 *Ascale, + syclex::fp8_e8m0 *Bscale, Tc *C, int M, int N, + int K, bool transpose_c = false, + bool colmajor_a = false, + bool colmajor_b = false, F &&lambda = {}) { + + for (unsigned int m = 0; m < M; m++) { + for (unsigned int n = 0; n < N; n++) { + int c_ind = transpose_c ? (n * M + m) : m * N + n; + Tc accb = *(C + c_ind); + // Per each block scale + for (unsigned int kb = 0; kb < K; kb += ScalingBlock) { + Tc acc = 0; + for (unsigned int k = 0; k < ScalingBlock; k++) { + + int a_ind = colmajor_a ? (kb * ScalingBlock + k * M + m) + : m * K + kb * ScalingBlock + k; + int b_ind = colmajor_b ? (n * K + kb * ScalingBlock + k) + : kb * ScalingBlock + k * N + n; + Ta *va = (Ta *)(A + a_ind * VF); + Tb *vb = (Tb *)(B + b_ind * VF); + + for (unsigned int i = 0; i < VF; i++) { + if constexpr (std::is_same_v && + std::is_same_v) + acc += make_fp32(va[i]) * make_fp32(vb[i]); + else if constexpr (std::is_same_v && + std::is_same_v) + acc += (float)va[i] * (float)vb[i]; + else if constexpr (std::is_same_v && + std::is_same_v || + std::is_integral_v && + std::is_integral_v || + (std::is_same_v || + std::is_same_v) || + (std::is_same_v && + std::is_same_v)) + acc += va[i] * vb[i]; + else + assert(false && "Unsupported type in matrix_multiply_ref."); + } + } // end k loop + // The 8-bit float types convert explicitly only, so the scales have to + // be cast before taking part in the multiplication. + accb += acc * (float)Ascale[m * K / ScalingBlock + kb] * + (float)Bscale[kb * N / ScalingBlock + n]; + } // end kb loop + + if constexpr (!std::is_same_v) { + lambda(accb); + } + *(C + c_ind) = accb; + } + } +} + template void matrix_apply(unsigned int rows, unsigned int cols, T *mat, F op) { for (unsigned int i = 0; i < rows; i++) diff --git a/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp index 9518d89e49d93..6aef13b503e7c 100644 --- a/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp +++ b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp @@ -6,6 +6,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// +#include template void assert_ops_ref(host_accessor mat, @@ -22,6 +23,21 @@ void assert_ops_ref(host_accessor mat, } } +// Overload for results held in a USM allocation instead of a buffer. +template +void assert_ops_ref(const T *mat, const float ref) { + for (size_t i = 0; i < NUM_ROWS; i++) + for (size_t j = 0; j < NUM_COLS; j++) { + float diff; + if constexpr (std::is_same_v) + diff = make_fp32(mat[i * NUM_COLS + j]) - ref; + else + diff = mat[i * NUM_COLS + j] - ref; + assert(std::fabs(static_cast(diff)) < + std::numeric_limits::epsilon()); + } +} + template @@ -65,6 +81,133 @@ void verify_op_ab(const T l, const T r, const float ref, OP op) { bufMat.get_host_access(read_only), ref); } +template , bool> = true> +void verify_op_ab(const sycl::half l, const sycl::half r, const float ref, + OP op) { + queue q; + size_t sg_size = get_sg_size(q); + + static constexpr size_t Rows = NUM_ROWS / VF; + static constexpr size_t Cols = NUM_COLS * VF; + + // The 8-bit float types only convert to half on the device, so the results + // are converted by matrix_copy's kernel rather than on the host. matrix_copy + // submits that kernel itself, so it must be called outside of any command + // group and both matrices have to live in USM. + T *mat = sycl::malloc_shared(Rows * Cols, q); + sycl::half *matH = sycl::malloc_shared(Rows * Cols, q); + + q.submit([&](handler &cgh) { + cgh.parallel_for( + nd_range<2>({NUM_ROWS / SUB_ROWS, NUM_COLS / SUB_COLS * sg_size}, + {1, 1 * sg_size}), + [=](nd_item<2> spmd_item) +#ifdef SG_SZ + [[sycl::reqd_sub_group_size(SG_SZ)]] +#endif + { + const auto global_idx = spmd_item.get_global_id(0); + const auto global_idy = spmd_item.get_global_id(1); + const auto sg_startx = global_idx - spmd_item.get_local_id(0); + const auto sg_starty = global_idy - spmd_item.get_local_id(1); + + sub_group sg = spmd_item.get_sub_group(); + joint_matrix sub_mat; + joint_matrix_fill(sg, sub_mat, T(l)); + joint_matrix_apply(sg, sub_mat, [=](T &x) { x = op((half)x, r); }); + auto pMat = + address_space_cast(mat); + ext::intel::experimental::matrix::joint_matrix_store( + sg, sub_mat, + pMat + (sg_startx * SUB_ROWS / VF) * Cols + + sg_starty / sg_size * SUB_COLS * VF, + Cols); + }); // parallel for + }).wait(); + + matrix_copy(q, Rows, Cols, mat, matH); + + assert_ops_ref(matH, ref); + + sycl::free(mat, q); + sycl::free(matH, q); +} + +template , bool> = true> +void verify_op_ab(const sycl::half l, const sycl::half r, const float ref, + OP op) { + queue q; + size_t sg_size = get_sg_size(q); + + // Rows/Cols count packed elements: each T holds numElems values, so a row of + // NUM_COLS values takes NUM_COLS / numElems of them, and the packed layout + // folds VF rows into one row of VF times the width. Cols is therefore the row + // stride in packed fp4_e2m1_x storage elements. + // NOTE: the packed VNNI layout for 4-bit types is unverified, GSD-9057. + static constexpr size_t Rows = NUM_ROWS / VF; + static constexpr size_t Cols = NUM_COLS / numElems * VF; + + // As in the fp8 overload above, the conversion to half happens in + // matrix_copy's own kernel, so both matrices live in USM. + T *mat = sycl::malloc_shared(Rows * Cols, q); + sycl::half *matH = sycl::malloc_shared(Rows * Cols * numElems, q); + + q.submit([&](handler &cgh) { + cgh.parallel_for( + nd_range<2>({NUM_ROWS / SUB_ROWS, NUM_COLS / SUB_COLS * sg_size}, + {1, 1 * sg_size}), + [=](nd_item<2> spmd_item) +#ifdef SG_SZ + [[sycl::reqd_sub_group_size(SG_SZ)]] +#endif + { + const auto global_idx = spmd_item.get_global_id(0); + const auto global_idy = spmd_item.get_global_id(1); + const auto sg_startx = global_idx - spmd_item.get_local_id(0); + const auto sg_starty = global_idy - spmd_item.get_local_id(1); + + sub_group sg = spmd_item.get_sub_group(); + joint_matrix sub_mat; + marray fillVal(l); + joint_matrix_fill(sg, sub_mat, T(fillVal)); + joint_matrix_apply(sg, sub_mat, [=](T &x) { + marray mval = + (marray)x; + for (unsigned int p = 0; p < numElems; p++) + mval[p] = op(mval[p], r); + // The 4-bit types construct from an marray explicitly only. + x = T(mval); + }); + auto pMat = + address_space_cast(mat); + // The row offset divides by VF, not by numElems: VF folds rows into + // the packed layout, while numElems is already accounted for by Cols + // being a count of packed elements. + ext::intel::experimental::matrix::joint_matrix_store( + sg, sub_mat, + pMat + (sg_startx * SUB_ROWS / VF) * Cols + + sg_starty / sg_size * SUB_COLS / numElems * VF, + Cols); + }); // parallel for + }).wait(); + + // matrix_copy expands each packed element into numElems halves, so it takes + // packed columns and writes numElems times as many values. + matrix_copy(q, Rows, Cols, mat, matH); + + assert_ops_ref(matH, ref); + + sycl::free(mat, q); + sycl::free(matH, q); +} + template void verify_op_c(const T l, const T r, const float ref, OP op) { @@ -252,6 +395,42 @@ int main() { break; } } + // fp4_e2m1_x packs 1 or 2 elements per byte, so the 4-bit tests run at a + // packing factor of 2. + constexpr unsigned int numElems = 2; + if (is_type_supported_by_device(q, matrix_type::fp4_e2m1)) { + // The advertised fp4_e2m1 combination is {msize=8, nsize=16, ksize=32} with + // no max_* sizes, so use::a and use::b extents have to match it exactly. + // + // VF counts logical elements, i.e. 32 bits / element bits: 2 for bfloat16, + // 4 for the 8-bit float types, 8 here. Eight 4-bit values are one 32-bit + // dword, so a dword holds VF consecutive k of a single B column. numElems + // is a separate axis: it only converts the logical column extent into a + // count of packed storage elements. This layout cannot be verified until + // IGC implements a 4-bit DPAS -- it currently rejects the i4 cooperative + // matrix component outright. Tracked by GSD-9057. + // + // If that instruction turns out to want the + // alternative format, where data is packed along the row first and whole + // bytes are folded afterwards, then VF becomes 4 and the host-side + // pack/fold order in joint_matrix_float4_impl.hpp has to change with it; + // the two are not independent. + test_ewops_ab, 8, 32, use::a, + layout::row_major, 1, sycl::half>(); + test_ewops_ab, 32, 16, use::b, + layout::ext_intel_packed, 8, sycl::half>(); + } + + if (is_type_supported_by_device(q, matrix_type::fp8_e5m2)) { + test_ewops_ab(); + test_ewops_ab(); + test_ewops_ab(); + test_ewops_ab(); + } return 0; } diff --git a/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp new file mode 100644 index 0000000000000..89d93ea642ec8 --- /dev/null +++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp @@ -0,0 +1,221 @@ +//==---------- joint_matrix_float4_impl.hpp - 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 +// +//===----------------------------------------------------------------------===// +#include +#include + +// The only advertised fp4_e2m1 combination is {msize=8, nsize=16, ksize=32}, +// so unlike the fp8 test TM is not a free parameter here. +constexpr size_t TM = 8; +constexpr size_t TN = 16; +constexpr size_t TK = 32; + +// numElems is the packing factor of fp4_e2m1_x: each storage element +// holds numElems 4-bit values. Matrix extents below are always expressed in +// logical (unpacked) elements, so offsets and strides into the packed A and B +// buffers divide by numElems. +// +// convertP selects where sycl::half is narrowed to fp4: when true, B is kept in +// memory as sycl::half and converted inside the kernel via +// joint_matrix_convert; when false, it is packed ahead of time by matrix_copy. +template +void joint_matrix_gemm_vnni( + sub_group sg, size_t sg_startx, size_t sg_starty, size_t sg_size, + multi_ptr + pA, + multi_ptr, + sycl::access::address_space::global_space, access::decorated::no> + pB, + multi_ptr + pC) { + joint_matrix sub_a; + joint_matrix sub_b; + joint_matrix sub_c; + + // B is addressed in packed storage elements unless it is still sycl::half. + constexpr size_t BPack = convertP ? 1 : numElems; + const size_t n_offset = sg_starty / sg_size * TN * vnniFactor; + + joint_matrix_load(sg, sub_c, + pC + (sg_startx * TM) * N + sg_starty / sg_size * TN, N, + layout::row_major); + for (int k = 0; k < K; k += TK) { + joint_matrix_load(sg, sub_a, pA + ((sg_startx * TM) * K + k) / numElems, + K / numElems); + if constexpr (convertP) { + joint_matrix sub_bh; + joint_matrix_load(sg, sub_bh, pB + (k * N + n_offset) / BPack, + N / BPack * vnniFactor); + joint_matrix_convert(sg, sub_bh, sub_b); + } else { + joint_matrix_load(sg, sub_b, pB + (k * N + n_offset) / BPack, + N / BPack * vnniFactor); + } + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + } + joint_matrix_store(sg, sub_c, + pC + (sg_startx * TM) * N + sg_starty / sg_size * TN, N, + layout::row_major); +} + +template +class fp4matrix; + +template +void matrix_multiply(TC *C, TA *A, + std::conditional_t *B, queue q) { + size_t NDRangeM = M / TM; + size_t NDRangeN = N / TN; + + auto pA = address_space_cast(A); + auto pB = address_space_cast(B); + auto pC = address_space_cast(C); + using kernel_name = fp4matrix; + size_t sg_size = get_sg_size(q); + q.submit([&](handler &cgh) { + cgh.parallel_for( + nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), + [=](nd_item<2> spmd_item) +#ifdef SG_SZ + [[sycl::reqd_sub_group_size(SG_SZ)]] +#endif + { + const auto global_idx = spmd_item.get_global_id(0); + const auto global_idy = spmd_item.get_global_id(1); + const auto sg_startx = global_idx - spmd_item.get_local_id(0); + const auto sg_starty = global_idy - spmd_item.get_local_id(1); + + sub_group sg = spmd_item.get_sub_group(); + joint_matrix_gemm_vnni(sg, sg_startx, sg_starty, + sg_size, pA, pB, pC); + }); // parallel for + }).wait(); +} + +// E2M1 only represents {0, 0.5, 1, 1.5, 2, 3, 4, 6} and their negations, so the +// random sycl::half input is not exactly representable. Round-tripping it +// through the 4-bit type makes the reference multiplication see exactly the +// values the device operates on. Unlike masking off fraction bits (what the fp8 +// test does via matrix_truncate_fraction_bits) this needs no knowledge of the +// bit layout. +template +void matrix_round_trip(queue q, unsigned int rows, unsigned int packedCols, + sycl::half *Mat, T4bit *Packed) { + matrix_copy(q, rows, packedCols, Mat, Packed); + matrix_copy(q, rows, packedCols, Packed, Mat); +} + +template +void joint_matrix_verify(queue q) { + sycl::half *Ah = malloc_shared(M * K, q); + sycl::half *Bh = malloc_shared(K * N, q); + TA *A = malloc_shared(M * K / numElems, q); + TB *B = malloc_shared(K * N / numElems, q); + TC *C = malloc_shared(M * N, q); + TC *D = malloc_shared(M * N, q); + + matrix_rand(M, K, Ah, 5); + matrix_rand(K, N, Bh, 5); + // Snap the reference data to values E2M1 represents exactly. + matrix_round_trip(q, M, K / numElems, Ah, A); + matrix_round_trip(q, K, N / numElems, Bh, B); + matrix_fill(M, N, C, (TC)1); + matrix_fill(M, N, D, (TC)1); + + if constexpr (vnniFactor > 1) { + // Apply VNNI on the sycl::half data, then pack it if the kernel expects + // fp4 in memory. + sycl::half *vnniBh = malloc_shared(K * N, q); + matrix_vnni(K, N, Bh, vnniBh, vnniFactor); + if constexpr (convertP) { + matrix_multiply(C, A, vnniBh, q); + } else { + TB *vnniB = malloc_shared(K * N / numElems, q); + matrix_copy(q, K, N / numElems, vnniBh, vnniB); + matrix_multiply(C, A, vnniB, q); + free(vnniB, q); + } + free(vnniBh, q); + } else { // row major + if constexpr (convertP) { + matrix_multiply(C, A, Bh, q); + } else { + matrix_multiply(C, A, B, q); + } + } + matrix_multiply_ref(Ah, Bh, D, M, N, K); + assert(matrix_compare(M, N, C, D)); + free(A, q); + free(B, q); + free(Ah, q); + free(Bh, q); + free(C, q); + free(D, q); +} + +template void fp4_combinations(queue q) { + // scale could be increased once these tests run on native hardware + static constexpr size_t SCALE = 2; + static constexpr size_t MATRIX_M = TM * SCALE; + // satisfy 64B stride requirement in 2D block load + static constexpr size_t MATRIX_N = std::max(TN * SCALE, 64ul); + static constexpr size_t MATRIX_K = TK * SCALE; + + using fp4 = syclex::fp4_e2m1_x; + + // vnniFactor 8 fills a 32-bit dword with 4-bit elements. joint_matrix_verify + // folds the unpacked sycl::half data by vnniFactor and only then packs pairs + // into bytes, so a dword ends up holding 8 consecutive k of one B column. + // Packing before folding would instead give a dword spanning two columns and + // four k, and vnniFactor would be 4; which of the two the hardware wants + // cannot be established until IGC implements a 4-bit DPAS. Tracked by + // GSD-9057. + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); +} + +int main() { + sycl::queue q; + if (!is_type_supported_by_device(q, matrix_type::fp4_e2m1)) { + std::cout << "fp4_e2m1 type not supported on this device" << std::endl; + return 0; + } + // A matrix element must pack two 4-bit values into each byte, so + // fp4_e2m1_x<2> is the only usable packing factor; fp4_e2m1_x<1> would leave + // the high nibble of every byte unused. + constexpr unsigned int numElems = 2; + fp4_combinations(q); +#if 0 + // Disabled by lack of bfloat16 accumulator support in IGC, as for the fp8 + // combinations in joint_matrix_float8_impl.hpp. Tracked by Jira GSD-10112. + fp4_combinations(q); +#endif + std::cout << "Passed\n"; + return 0; +} diff --git a/sycl/test-e2e/Matrix/Inputs/joint_matrix_float8_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float8_impl.hpp new file mode 100644 index 0000000000000..275d7d7a867d6 --- /dev/null +++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float8_impl.hpp @@ -0,0 +1,185 @@ +//==---------- joint_matrix_float8_impl.hpp - 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 +// +//===----------------------------------------------------------------------===// +#include +#include + +constexpr size_t TN = 16; +constexpr size_t TK = 32; + +template +void joint_matrix_gemm_vnni( + sub_group sg, size_t sg_startx, size_t sg_starty, size_t sg_size, + multi_ptr + pA, + multi_ptr + pB, + multi_ptr + pC) { + joint_matrix sub_a; + joint_matrix sub_b; + joint_matrix sub_c; + joint_matrix_load(sg, sub_c, + pC + (sg_startx * TM) * N + sg_starty / sg_size * TN, N, + layout::row_major); + for (int k = 0; k < K; k += TK) { + joint_matrix_load(sg, sub_a, pA + (sg_startx * TM) * K + k, K); + joint_matrix_load(sg, sub_b, + pB + k * N + sg_starty / sg_size * TN * vnniFactor, + N * vnniFactor); + joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c); + } + joint_matrix_store(sg, sub_c, + pC + (sg_startx * TM) * N + sg_starty / sg_size * TN, N, + layout::row_major); +} + +template +class fp8matrix; + +template +void matrix_multiply(TC *C, TA *A, TB *B, queue q) { + size_t NDRangeM = M / TM; + size_t NDRangeN = N / TN; + + auto pA = address_space_cast(A); + auto pB = address_space_cast(B); + auto pC = address_space_cast(C); + size_t sg_size = + get_sg_size>(q); + q.submit([&](handler &cgh) { + cgh.parallel_for>( + nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), + [=](nd_item<2> spmd_item) +#ifdef SG_SZ + [[sycl::reqd_sub_group_size(SG_SZ)]] +#endif + { + const auto global_idx = spmd_item.get_global_id(0); + const auto global_idy = spmd_item.get_global_id(1); + const auto sg_startx = global_idx - spmd_item.get_local_id(0); + const auto sg_starty = global_idy - spmd_item.get_local_id(1); + + sub_group sg = spmd_item.get_sub_group(); + joint_matrix_gemm_vnni(sg, sg_startx, sg_starty, sg_size, + pA, pB, pC); + }); // parallel for + }).wait(); +} + +template +void joint_matrix_verify(queue q) { + sycl::half *Ah = malloc_shared(M * K, q); + sycl::half *Bh = malloc_shared(K * N, q); + TA *A = malloc_shared(M * K, q); + TB *B = malloc_shared(K * N, q); + TC *C = malloc_shared(M * N, q); + TC *D = malloc_shared(M * N, q); + + matrix_rand(M, K, Ah, 5); + matrix_truncate_fraction_bits(M, K, Ah); + matrix_rand(K, N, Bh, 5); + matrix_truncate_fraction_bits(K, N, Bh); + matrix_fill(M, N, C, (TC)1); + matrix_fill(M, N, D, (TC)1); + // Assign Ah and Bh values to A and B + matrix_copy(q, M, K, Ah, A); + matrix_copy(q, K, N, Bh, B); + + if (vnniFactor > 1) { + TB *vnniB = malloc_shared(K * N, q); + matrix_vnni(K, N, B, vnniB, vnniFactor); + matrix_multiply(C, A, vnniB, q); + } else { + matrix_multiply( + C, A, B, q); + } + matrix_multiply_ref(Ah, Bh, D, M, N, K); + assert(matrix_compare(M, N, C, D)); + free(A, q); + free(B, q); + free(Ah, q); + free(Bh, q); + free(C, q); + free(D, q); +} + +template void bf8_hf8_combinations(queue q) { + // scale could be increased to 8 once these tests run on native hardware + static constexpr size_t SCALE = 2; + static constexpr size_t MATRIX_M = TM * SCALE; + // satisfy 64B stride requirement in 2D block load + static constexpr size_t MATRIX_N = std::max(TN * SCALE, 64ul); + static constexpr size_t MATRIX_K = TK * SCALE; + + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); +#if 0 + // These combinations are disabled by lack of bfloat16 accumulator support in IGC + // Adding these is tracked by Jira GSD-10112 + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); + joint_matrix_verify(q); +#endif +} + +int main() { + sycl::queue q; + if (!is_type_supported_by_device(q, matrix_type::fp8_e5m2)) { + std::cout << "bf8 and hf8 types not supported on this device" << std::endl; + return 0; + } + bf8_hf8_combinations<1 /*TM*/>(q); + bf8_hf8_combinations<2 /*TM*/>(q); + bf8_hf8_combinations<3 /*TM*/>(q); + bf8_hf8_combinations<4 /*TM*/>(q); + bf8_hf8_combinations<5 /*TM*/>(q); + bf8_hf8_combinations<6 /*TM*/>(q); + bf8_hf8_combinations<7 /*TM*/>(q); + bf8_hf8_combinations<8 /*TM*/>(q); + std::cout << "Passed\n"; + return 0; +} diff --git a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops.cpp b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops.cpp index 2e2a1b88c9298..5e2d55d0d0603 100644 --- a/sycl/test-e2e/Matrix/SG32/element_wise_all_ops.cpp +++ b/sycl/test-e2e/Matrix/SG32/element_wise_all_ops.cpp @@ -13,7 +13,7 @@ // UNSUPPORTED: gpu-intel-dg2 // UNSUPPORTED-TRACKER: GSD-10700 -// RUN: %{build} -o %t.out +// RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_EXT_float8,+SPV_INTEL_float4,+SPV_INTEL_int4,+SPV_INTEL_fp_conversions -o %t.out // RUN: %{run} %t.out #include "common.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_float4.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_float4.cpp new file mode 100644 index 0000000000000..f8a1d25df2387 --- /dev/null +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_float4.cpp @@ -0,0 +1,19 @@ +//==----------- joint_matrix_float4.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: arch-intel_gpu_cri, aspect-ext_intel_matrix + +// RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_INTEL_fp_conversions,+SPV_INTEL_float4,+SPV_INTEL_int4,+SPV_KHR_bfloat16 -o %t.out +// RUN: %{run} %t.out + +// UNSUPPORTED: target-nvidia, target-amd, spirv-backend +// UNSUPPORTED-INTENDED: only supported by backends with CRI driver, and the +// SPIR-V backend does not support the required SPIR-V extensions + +#include "common.hpp" +#define SG_SZ 32 +#include "joint_matrix_float4_impl.hpp" diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_float8.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_float8.cpp new file mode 100644 index 0000000000000..5a015a64f3ccf --- /dev/null +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_float8.cpp @@ -0,0 +1,15 @@ +//==----------- joint_matrix_float8.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: arch-intel_gpu_cri, aspect-ext_intel_matrix + +// RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_EXT_float8,+SPV_INTEL_fp_conversions,+SPV_KHR_bfloat16 -o %t.out +// RUN: %{run} %t.out + +#include "common.hpp" +#define SG_SZ 32 +#include "joint_matrix_float8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops.cpp b/sycl/test-e2e/Matrix/element_wise_all_ops.cpp index 3fb8786b028f5..cb8ae9a016bf7 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops.cpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops.cpp @@ -9,7 +9,7 @@ // REQUIRES: aspect-ext_intel_matrix -// RUN: %{build} -o %t.out +// RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_EXT_float8,+SPV_INTEL_float4,+SPV_INTEL_int4,+SPV_INTEL_fp_conversions -o %t.out // RUN: %{run} %t.out #include "common.hpp" diff --git a/sycl/test-e2e/Matrix/joint_matrix_float4.cpp b/sycl/test-e2e/Matrix/joint_matrix_float4.cpp new file mode 100644 index 0000000000000..64784a5f73750 --- /dev/null +++ b/sycl/test-e2e/Matrix/joint_matrix_float4.cpp @@ -0,0 +1,18 @@ +//==----------- joint_matrix_float4.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: arch-intel_gpu_cri, aspect-ext_intel_matrix + +// RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_INTEL_fp_conversions,+SPV_INTEL_float4,+SPV_INTEL_int4,+SPV_KHR_bfloat16 -o %t.out +// RUN: %{run} %t.out + +// UNSUPPORTED: target-nvidia, target-amd, spirv-backend +// UNSUPPORTED-INTENDED: only supported by backends with CRI driver, and the +// SPIR-V backend does not support the required SPIR-V extensions + +#include "common.hpp" +#include "joint_matrix_float4_impl.hpp" diff --git a/sycl/test-e2e/Matrix/joint_matrix_float8.cpp b/sycl/test-e2e/Matrix/joint_matrix_float8.cpp new file mode 100644 index 0000000000000..33082e299be85 --- /dev/null +++ b/sycl/test-e2e/Matrix/joint_matrix_float8.cpp @@ -0,0 +1,14 @@ +//==----------- joint_matrix_float8.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: arch-intel_gpu_cri, aspect-ext_intel_matrix + +// RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_EXT_float8,+SPV_INTEL_fp_conversions,+SPV_KHR_bfloat16 -o %t.out +// RUN: %{run} %t.out + +#include "common.hpp" +#include "joint_matrix_float8_impl.hpp" From 59131b4624ffd61bc70a9c027382570c0e833fe4 Mon Sep 17 00:00:00 2001 From: Dounia Khaldi Date: Mon, 31 Aug 2026 19:09:38 +0000 Subject: [PATCH 2/6] Cleanup --- clang/lib/CodeGen/CodeGenTypes.cpp | 2 - .../oneapi/matrix/matrix-unified-utils.hpp | 6 +- sycl/include/sycl/sycl.hpp | 4 +- sycl/source/detail/device_impl.hpp | 3 +- sycl/test-e2e/Matrix/Inputs/common.hpp | 11 +-- .../Inputs/element_wise_all_ops_impl.hpp | 94 ++++++------------- .../Inputs/joint_matrix_float4_impl.hpp | 10 +- 7 files changed, 40 insertions(+), 90 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index da81390598ea7..8b19136688970 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -399,8 +399,6 @@ llvm::Type *CodeGenTypes::ConvertSPVCooperativeMatrixType(RecordDecl *RD) { CompTy = llvm::Type::getFloatTy(getLLVMContext()); } else if (LlvmTyName == "bfloat16") { CompTy = llvm::Type::getInt16Ty(getLLVMContext()); - // The 8-bit float types are class templates parameterized by the packing - // width (fp8_e5m2_x), so match on the class name without the width. } else if (LlvmTyName == "fp8_e5m2_x" || LlvmTyName == "fp8_e4m3_x" || LlvmTyName == "fp8_e8m0_x") { CompTy = llvm::Type::getInt8Ty(getLLVMContext()); diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp index 21fc4daf4b7e5..efbf55bb98bba 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp @@ -12,8 +12,8 @@ #include // std::string_view #include // __spv namespace #include // bfloat16 -#include // fp4_e2m1_x -#include // fp8_e5m2, fp8_e4m3 +#include // for fp4_e2m1_x +#include // for fp8_e5m2 #include // std::pair namespace sycl { @@ -88,8 +88,6 @@ extern "C" constexpr __spv::MatrixLayout joint_matrix_layout_to_spv( } } -// fp4_e2m1_x is a packed type: any supported packing width N designates the -// same E2M1 element format, so matrix queries match on the format, not on N. template struct is_fp4_e2m1 : std::false_type {}; template struct is_fp4_e2m1> diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index a4cb9f99481e4..237d3e606723b 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -125,6 +125,8 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.") #include #include #include +#include +#include #include #include #include @@ -150,8 +152,6 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.") #include #include #include -#include -#include #include #include #include diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index 99a61ee1a3c0c..d1588af1eb243 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -2207,7 +2207,8 @@ class device_impl { matrix_type::fp32, matrix_type::fp32}, }; // fp8 and fp4_e2m1 are supported starting with Crescent Island. - if (architecture::intel_gpu_cri == DeviceArch) + bool SupportsFP8AndFP4 = (architecture::intel_gpu_cri == DeviceArch); + if (SupportsFP8AndFP4) pvc_combs.insert( pvc_combs.end(), { diff --git a/sycl/test-e2e/Matrix/Inputs/common.hpp b/sycl/test-e2e/Matrix/Inputs/common.hpp index dbe42a5f85f4c..b870a4d591b4d 100644 --- a/sycl/test-e2e/Matrix/Inputs/common.hpp +++ b/sycl/test-e2e/Matrix/Inputs/common.hpp @@ -179,9 +179,9 @@ void matrix_copy(unsigned int rows, unsigned int cols, T *src, T *dst) { // 8-bit float types: they pack a single element and convert to/from half. template -constexpr bool is_fp8_type_v = std::is_same_v || - std::is_same_v || - std::is_same_v; +constexpr bool is_fp8_type_v = + std::is_same_v || + std::is_same_v || std::is_same_v; // 4-bit types: they pack numElems elements and convert to/from // marray. @@ -189,8 +189,6 @@ template constexpr bool is_4bit_type_v = std::is_same_v>; // numElems is the packing factor for the 4bits types. It can be 2 or 8 -// src is const so that Ts is deduced without a cv-qualifier when it comes from -// a read-only accessor: the is_*_type_v checks below match unqualified types. template void matrix_copy(queue q, unsigned int rows, unsigned int cols, const Ts *src, Td *dst) { @@ -199,8 +197,7 @@ void matrix_copy(queue q, unsigned int rows, unsigned int cols, const Ts *src, for (unsigned int j = 0; j < cols; j++) { if constexpr (std::is_same_v && is_fp8_type_v) dst[i * cols + j] = (Td)src[i * cols + j]; - else if constexpr (std::is_same_v && - is_fp8_type_v) + else if constexpr (std::is_same_v && is_fp8_type_v) dst[i * cols + j] = src[i * cols + j]; else if constexpr (std::is_same_v && is_4bit_type_v) { diff --git a/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp index 6aef13b503e7c..89b5be563f12d 100644 --- a/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp +++ b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp @@ -1,4 +1,5 @@ #include +#include //==----------- element_wise_all_ops_impl.hpp - DPC++ joint_matrix---------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. @@ -6,7 +7,6 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -#include template void assert_ops_ref(host_accessor mat, @@ -47,7 +47,6 @@ void verify_op_ab(const T l, const T r, const float ref, OP op) { buffer bufMat(big_mat.get_data(), range<2>(NUM_ROWS / VF, NUM_COLS * VF)); - queue q; size_t sg_size = get_sg_size(q); q.submit([&](handler &cgh) { @@ -87,18 +86,10 @@ template (NUM_ROWS / VF * NUM_COLS * VF, q); + sycl::half *matH = + malloc_shared(NUM_ROWS / VF * NUM_COLS * VF, q); size_t sg_size = get_sg_size(q); - - static constexpr size_t Rows = NUM_ROWS / VF; - static constexpr size_t Cols = NUM_COLS * VF; - - // The 8-bit float types only convert to half on the device, so the results - // are converted by matrix_copy's kernel rather than on the host. matrix_copy - // submits that kernel itself, so it must be called outside of any command - // group and both matrices have to live in USM. - T *mat = sycl::malloc_shared(Rows * Cols, q); - sycl::half *matH = sycl::malloc_shared(Rows * Cols, q); - q.submit([&](handler &cgh) { cgh.parallel_for( nd_range<2>({NUM_ROWS / SUB_ROWS, NUM_COLS / SUB_COLS * sg_size}, @@ -113,27 +104,28 @@ void verify_op_ab(const sycl::half l, const sycl::half r, const float ref, const auto sg_startx = global_idx - spmd_item.get_local_id(0); const auto sg_starty = global_idy - spmd_item.get_local_id(1); + auto pMat = + address_space_cast(mat); + sub_group sg = spmd_item.get_sub_group(); joint_matrix sub_mat; joint_matrix_fill(sg, sub_mat, T(l)); joint_matrix_apply(sg, sub_mat, [=](T &x) { x = op((half)x, r); }); - auto pMat = - address_space_cast(mat); ext::intel::experimental::matrix::joint_matrix_store( sg, sub_mat, - pMat + (sg_startx * SUB_ROWS / VF) * Cols + + pMat + (sg_startx * SUB_ROWS / VF) * NUM_COLS * VF + sg_starty / sg_size * SUB_COLS * VF, - Cols); + NUM_COLS * VF); }); // parallel for }).wait(); - matrix_copy(q, Rows, Cols, mat, matH); + matrix_copy(q, NUM_ROWS, NUM_COLS, mat, matH); - assert_ops_ref(matH, ref); + assert_ops_ref(matH, ref); - sycl::free(mat, q); - sycl::free(matH, q); + free(mat, q); + free(matH, q); } template , bool> = true> void verify_op_ab(const sycl::half l, const sycl::half r, const float ref, OP op) { + // Cols is the row stride in packed fp4_e2m1_x storage elements. + // NOTE: the packed VNNI layout for 4-bit types is unverified GSD-9057 + constexpr size_t Cols = NUM_COLS / numElems * VF; + queue q; + T *mat = malloc_shared(NUM_ROWS / VF * Cols, q); + sycl::half *matH = + malloc_shared(NUM_ROWS / VF * NUM_COLS * VF, q); size_t sg_size = get_sg_size(q); - - // Rows/Cols count packed elements: each T holds numElems values, so a row of - // NUM_COLS values takes NUM_COLS / numElems of them, and the packed layout - // folds VF rows into one row of VF times the width. Cols is therefore the row - // stride in packed fp4_e2m1_x storage elements. - // NOTE: the packed VNNI layout for 4-bit types is unverified, GSD-9057. - static constexpr size_t Rows = NUM_ROWS / VF; - static constexpr size_t Cols = NUM_COLS / numElems * VF; - - // As in the fp8 overload above, the conversion to half happens in - // matrix_copy's own kernel, so both matrices live in USM. - T *mat = sycl::malloc_shared(Rows * Cols, q); - sycl::half *matH = sycl::malloc_shared(Rows * Cols * numElems, q); - q.submit([&](handler &cgh) { cgh.parallel_for( nd_range<2>({NUM_ROWS / SUB_ROWS, NUM_COLS / SUB_COLS * sg_size}, @@ -172,6 +157,10 @@ void verify_op_ab(const sycl::half l, const sycl::half r, const float ref, const auto sg_startx = global_idx - spmd_item.get_local_id(0); const auto sg_starty = global_idy - spmd_item.get_local_id(1); + auto pMat = + address_space_cast(mat); + sub_group sg = spmd_item.get_sub_group(); joint_matrix sub_mat; marray fillVal(l); @@ -184,12 +173,6 @@ void verify_op_ab(const sycl::half l, const sycl::half r, const float ref, // The 4-bit types construct from an marray explicitly only. x = T(mval); }); - auto pMat = - address_space_cast(mat); - // The row offset divides by VF, not by numElems: VF folds rows into - // the packed layout, while numElems is already accounted for by Cols - // being a count of packed elements. ext::intel::experimental::matrix::joint_matrix_store( sg, sub_mat, pMat + (sg_startx * SUB_ROWS / VF) * Cols + @@ -198,14 +181,12 @@ void verify_op_ab(const sycl::half l, const sycl::half r, const float ref, }); // parallel for }).wait(); - // matrix_copy expands each packed element into numElems halves, so it takes - // packed columns and writes numElems times as many values. - matrix_copy(q, Rows, Cols, mat, matH); + matrix_copy(q, NUM_ROWS / VF, Cols, mat, matH); - assert_ops_ref(matH, ref); + assert_ops_ref(matH, ref); - sycl::free(mat, q); - sycl::free(matH, q); + free(mat, q); + free(matH, q); } template >( Tv(5.0), Tv(2.0), 7.0, [](auto l, auto r) { return l + r; }); @@ -399,22 +379,6 @@ int main() { // packing factor of 2. constexpr unsigned int numElems = 2; if (is_type_supported_by_device(q, matrix_type::fp4_e2m1)) { - // The advertised fp4_e2m1 combination is {msize=8, nsize=16, ksize=32} with - // no max_* sizes, so use::a and use::b extents have to match it exactly. - // - // VF counts logical elements, i.e. 32 bits / element bits: 2 for bfloat16, - // 4 for the 8-bit float types, 8 here. Eight 4-bit values are one 32-bit - // dword, so a dword holds VF consecutive k of a single B column. numElems - // is a separate axis: it only converts the logical column extent into a - // count of packed storage elements. This layout cannot be verified until - // IGC implements a 4-bit DPAS -- it currently rejects the i4 cooperative - // matrix component outright. Tracked by GSD-9057. - // - // If that instruction turns out to want the - // alternative format, where data is packed along the row first and whole - // bytes are folded afterwards, then VF becomes 4 and the host-side - // pack/fold order in joint_matrix_float4_impl.hpp has to change with it; - // the two are not independent. test_ewops_ab, 8, 32, use::a, layout::row_major, 1, sycl::half>(); test_ewops_ab, 32, 16, use::b, diff --git a/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp index 89d93ea642ec8..3421b6fc47b92 100644 --- a/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp +++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp @@ -8,8 +8,6 @@ #include #include -// The only advertised fp4_e2m1 combination is {msize=8, nsize=16, ksize=32}, -// so unlike the fp8 test TM is not a free parameter here. constexpr size_t TM = 8; constexpr size_t TN = 16; constexpr size_t TK = 32; @@ -183,13 +181,7 @@ template void fp4_combinations(queue q) { using fp4 = syclex::fp4_e2m1_x; - // vnniFactor 8 fills a 32-bit dword with 4-bit elements. joint_matrix_verify - // folds the unpacked sycl::half data by vnniFactor and only then packs pairs - // into bytes, so a dword ends up holding 8 consecutive k of one B column. - // Packing before folding would instead give a dword spanning two columns and - // four k, and vnniFactor would be 4; which of the two the hardware wants - // cannot be established until IGC implements a 4-bit DPAS. Tracked by - // GSD-9057. + // vnniFactor 8 fills a 32-bit dword with 4-bit elements joint_matrix_verify(q); joint_matrix_verify Date: Tue, 1 Sep 2026 13:56:33 +0000 Subject: [PATCH 3/6] Add missing fp8, fp4 flags on other tests that touch them --- sycl/test-e2e/Matrix/element_wise_all_ops_1d.cpp | 2 +- sycl/test-e2e/Matrix/element_wise_all_ops_1d_cont.cpp | 2 +- sycl/test-e2e/Matrix/element_wise_all_ops_scalar.cpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_1d.cpp b/sycl/test-e2e/Matrix/element_wise_all_ops_1d.cpp index 89bb84b33acea..7f7da370fc16d 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_1d.cpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_1d.cpp @@ -13,7 +13,7 @@ // XFAIL: windows && intel_gpu_lnl_m && O0 // XFAIL-TRACKER: CMPLRLLVM-72111 -// RUN: %{build} -o %t.out +// RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_EXT_float8,+SPV_INTEL_float4,+SPV_INTEL_int4,+SPV_INTEL_fp_conversions -o %t.out // RUN: env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out #include "common.hpp" diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_1d_cont.cpp b/sycl/test-e2e/Matrix/element_wise_all_ops_1d_cont.cpp index 4806c1fadcebf..d17f185a3f7b5 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_1d_cont.cpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_1d_cont.cpp @@ -9,7 +9,7 @@ // REQUIRES: aspect-ext_intel_matrix, gpu // REQUIRES-INTEL-DRIVER: lin: 30049 -// RUN: %{build} -o %t.out +// RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_EXT_float8,+SPV_INTEL_float4,+SPV_INTEL_int4,+SPV_INTEL_fp_conversions -o %t.out // RUN: env IGC_JointMatrixLoadStoreOpt=2 %{run} %t.out #include "common.hpp" diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_scalar.cpp b/sycl/test-e2e/Matrix/element_wise_all_ops_scalar.cpp index 04702b066102a..4d38ece14de24 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_scalar.cpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_scalar.cpp @@ -9,7 +9,7 @@ // REQUIRES: aspect-ext_intel_matrix, gpu // REQUIRES-INTEL-DRIVER: lin: 30049 -// RUN: %{build} -o %t.out +// RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_EXT_float8,+SPV_INTEL_float4,+SPV_INTEL_int4,+SPV_INTEL_fp_conversions -o %t.out // RUN: env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out #include "common.hpp" From 1300888f7b8844d0316334cac9870cbbc9a8251e Mon Sep 17 00:00:00 2001 From: Dounia Khaldi Date: Tue, 1 Sep 2026 16:08:47 +0000 Subject: [PATCH 4/6] add xfails --- sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp | 4 ++++ sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp | 3 +-- sycl/test-e2e/Matrix/SG32/joint_matrix_float4.cpp | 3 +++ sycl/test-e2e/Matrix/SG32/joint_matrix_float8.cpp | 4 ++++ sycl/test-e2e/Matrix/joint_matrix_float4.cpp | 3 +++ sycl/test-e2e/Matrix/joint_matrix_float8.cpp | 4 ++++ 6 files changed, 19 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp index 89b5be563f12d..8ae3a2283035e 100644 --- a/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp +++ b/sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp @@ -375,6 +375,9 @@ int main() { break; } } +#if 0 + // Disabled by lack of 4-bit DPAS support in IGC; the packed VNNI layout for + // the 4-bit types is unverified as a result. Tracked by Jira GSD-9057. // fp4_e2m1_x packs 1 or 2 elements per byte, so the 4-bit tests run at a // packing factor of 2. constexpr unsigned int numElems = 2; @@ -384,6 +387,7 @@ int main() { test_ewops_ab, 32, 16, use::b, layout::ext_intel_packed, 8, sycl::half>(); } +#endif if (is_type_supported_by_device(q, matrix_type::fp8_e5m2)) { test_ewops_ab is the only usable packing factor; fp4_e2m1_x<1> would leave - // the high nibble of every byte unused. + // fp4_e2m1_x<2> is the only usable packing factor constexpr unsigned int numElems = 2; fp4_combinations(q); #if 0 diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_float4.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_float4.cpp index f8a1d25df2387..0bff8c613baa6 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_float4.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_float4.cpp @@ -7,6 +7,9 @@ //===----------------------------------------------------------------------===// // REQUIRES: arch-intel_gpu_cri, aspect-ext_intel_matrix +// XFAIL: arch-intel_gpu_cri +// XFAIL-TRACKER: GSD-9057 + // RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_INTEL_fp_conversions,+SPV_INTEL_float4,+SPV_INTEL_int4,+SPV_KHR_bfloat16 -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/SG32/joint_matrix_float8.cpp b/sycl/test-e2e/Matrix/SG32/joint_matrix_float8.cpp index 5a015a64f3ccf..cac8b48eca9bf 100644 --- a/sycl/test-e2e/Matrix/SG32/joint_matrix_float8.cpp +++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_float8.cpp @@ -10,6 +10,10 @@ // RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_EXT_float8,+SPV_INTEL_fp_conversions,+SPV_KHR_bfloat16 -o %t.out // RUN: %{run} %t.out +// UNSUPPORTED: target-nvidia, target-amd, spirv-backend +// UNSUPPORTED-INTENDED: only supported by backends with CRI driver, and the +// SPIR-V backend does not support the required SPIR-V extensions + #include "common.hpp" #define SG_SZ 32 #include "joint_matrix_float8_impl.hpp" diff --git a/sycl/test-e2e/Matrix/joint_matrix_float4.cpp b/sycl/test-e2e/Matrix/joint_matrix_float4.cpp index 64784a5f73750..7403f31b34e9d 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_float4.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_float4.cpp @@ -7,6 +7,9 @@ //===----------------------------------------------------------------------===// // REQUIRES: arch-intel_gpu_cri, aspect-ext_intel_matrix +// XFAIL: arch-intel_gpu_cri +// XFAIL-TRACKER: GSD-9057 + // RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_INTEL_fp_conversions,+SPV_INTEL_float4,+SPV_INTEL_int4,+SPV_KHR_bfloat16 -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/Matrix/joint_matrix_float8.cpp b/sycl/test-e2e/Matrix/joint_matrix_float8.cpp index 33082e299be85..59824d20aa06f 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_float8.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_float8.cpp @@ -10,5 +10,9 @@ // RUN: %{build} -Xspirv-translator=spir64 --spirv-ext=+SPV_EXT_float8,+SPV_INTEL_fp_conversions,+SPV_KHR_bfloat16 -o %t.out // RUN: %{run} %t.out +// UNSUPPORTED: target-nvidia, target-amd, spirv-backend +// UNSUPPORTED-INTENDED: only supported by backends with CRI driver, and the +// SPIR-V backend does not support the required SPIR-V extensions + #include "common.hpp" #include "joint_matrix_float8_impl.hpp" From 0b411093986137b21da8286b366f61764c18e832 Mon Sep 17 00:00:00 2001 From: Dounia Khaldi Date: Tue, 1 Sep 2026 19:41:00 +0000 Subject: [PATCH 5/6] Removed unecesary includes and fix a bug in the fp4 header --- .../sycl/ext/oneapi/experimental/float_4bit/types.hpp | 8 +++++--- sycl/include/sycl/sycl.hpp | 2 -- sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp | 2 +- sycl/test-e2e/Matrix/Inputs/joint_matrix_float8_impl.hpp | 2 +- 4 files changed, 7 insertions(+), 7 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/float_4bit/types.hpp b/sycl/include/sycl/ext/oneapi/experimental/float_4bit/types.hpp index fef63bd8af6a0..f3e029df7c5b0 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/float_4bit/types.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/float_4bit/types.hpp @@ -30,12 +30,14 @@ #ifdef __SYCL_DEVICE_ONLY__ namespace sycl { +inline namespace _V1 { namespace detail { using fp4_float16_vec2 = _Float16 __attribute__((ext_vector_type(2))); using fp4_bfloat16_vec2 = __bf16 __attribute__((ext_vector_type(2))); using fp4_uint8_vec2 = uint8_t __attribute__((ext_vector_type(2))); using fp4_uint8_vec1 = uint8_t __attribute__((ext_vector_type(1))); } // namespace detail +} // namespace _V1 } // namespace sycl // FP4 builtins. The SPIR-V translator maps these to SPV_INTEL_float4 / @@ -392,10 +394,10 @@ ConvertFromFP4ToBinaryFloat_CPU(uint8_t code, if (BitWidth(magnitude) > Traits::ValueBits) { if constexpr (Traits::IsSigned) - return negative ? std::numeric_limits::min() - : std::numeric_limits::max(); + return negative ? (std::numeric_limits::min)() + : (std::numeric_limits::max)(); else - return negative ? ToT{0} : std::numeric_limits::max(); + return negative ? ToT{0} : (std::numeric_limits::max)(); } const UnsignedT narrowed = static_cast(magnitude); diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index 237d3e606723b..d133824184470 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -125,8 +125,6 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.") #include #include #include -#include -#include #include #include #include diff --git a/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp index 8a14efc2aaac1..015103a97728c 100644 --- a/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp +++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp @@ -176,7 +176,7 @@ template void fp4_combinations(queue q) { static constexpr size_t SCALE = 2; static constexpr size_t MATRIX_M = TM * SCALE; // satisfy 64B stride requirement in 2D block load - static constexpr size_t MATRIX_N = std::max(TN * SCALE, 64ul); + static constexpr size_t MATRIX_N = std::max(TN * SCALE, 64); static constexpr size_t MATRIX_K = TK * SCALE; using fp4 = syclex::fp4_e2m1_x; diff --git a/sycl/test-e2e/Matrix/Inputs/joint_matrix_float8_impl.hpp b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float8_impl.hpp index 275d7d7a867d6..4a2d5eab71f79 100644 --- a/sycl/test-e2e/Matrix/Inputs/joint_matrix_float8_impl.hpp +++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float8_impl.hpp @@ -125,7 +125,7 @@ template void bf8_hf8_combinations(queue q) { static constexpr size_t SCALE = 2; static constexpr size_t MATRIX_M = TM * SCALE; // satisfy 64B stride requirement in 2D block load - static constexpr size_t MATRIX_N = std::max(TN * SCALE, 64ul); + static constexpr size_t MATRIX_N = std::max(TN * SCALE, 64); static constexpr size_t MATRIX_K = TK * SCALE; joint_matrix_verify Date: Tue, 1 Sep 2026 22:11:05 +0000 Subject: [PATCH 6/6] forward declare the fp4/fp8 types to avoid includes in the headers --- .../oneapi/experimental/float_4bit/types.hpp | 8 ++- .../oneapi/matrix/matrix-unified-utils.hpp | 53 +++++++++++-------- .../sycl/ext/oneapi/matrix/matrix-unified.hpp | 14 +++-- .../sycl/ext/oneapi/matrix/query-types.hpp | 8 +-- sycl/test-e2e/Matrix/Inputs/common.hpp | 4 ++ 5 files changed, 47 insertions(+), 40 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/float_4bit/types.hpp b/sycl/include/sycl/ext/oneapi/experimental/float_4bit/types.hpp index f3e029df7c5b0..fef63bd8af6a0 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/float_4bit/types.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/float_4bit/types.hpp @@ -30,14 +30,12 @@ #ifdef __SYCL_DEVICE_ONLY__ namespace sycl { -inline namespace _V1 { namespace detail { using fp4_float16_vec2 = _Float16 __attribute__((ext_vector_type(2))); using fp4_bfloat16_vec2 = __bf16 __attribute__((ext_vector_type(2))); using fp4_uint8_vec2 = uint8_t __attribute__((ext_vector_type(2))); using fp4_uint8_vec1 = uint8_t __attribute__((ext_vector_type(1))); } // namespace detail -} // namespace _V1 } // namespace sycl // FP4 builtins. The SPIR-V translator maps these to SPV_INTEL_float4 / @@ -394,10 +392,10 @@ ConvertFromFP4ToBinaryFloat_CPU(uint8_t code, if (BitWidth(magnitude) > Traits::ValueBits) { if constexpr (Traits::IsSigned) - return negative ? (std::numeric_limits::min)() - : (std::numeric_limits::max)(); + return negative ? std::numeric_limits::min() + : std::numeric_limits::max(); else - return negative ? ToT{0} : (std::numeric_limits::max)(); + return negative ? ToT{0} : std::numeric_limits::max(); } const UnsignedT narrowed = static_cast(magnitude); diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp index efbf55bb98bba..965fa6d9cfdd6 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp @@ -8,19 +8,30 @@ #pragma once +#include // size_t #include // std::optional #include // std::string_view #include // __spv namespace #include // bfloat16 -#include // for fp4_e2m1_x -#include // for fp8_e5m2 -#include // std::pair +#include // std::pair namespace sycl { inline namespace _V1 { namespace ext { namespace oneapi { namespace experimental { + +// The matrix interfaces only ever name the 4-bit and 8-bit floating point types +// as template arguments, so forward declarations are enough here. Including +// and its 8-bit counterpart +// instead would pull both extensions into the transitive closure of +// , which they are deliberately not part of: code that uses +// fp4/fp8 includes those headers explicitly. +template class fp4_e2m1_x; +template class fp8_e4m3_x; +template class fp8_e5m2_x; +template class fp8_e8m0_x; + namespace matrix { enum class use { a, b, accumulator, scale }; @@ -112,38 +123,34 @@ constexpr uint32_t CalculateMatrixOperand() { if constexpr (std::is_signed::value) returnValue += static_cast( __spv::MatrixOperands::MatrixBSignedComponentsKHR); - if constexpr (std::is_same< - Ta, sycl::ext::oneapi::experimental::fp8_e5m2>::value && - std::is_same< - Tb, sycl::ext::oneapi::experimental::fp8_e5m2>::value && - std::is_same::value) + if constexpr ( + std::is_same>::value && + std::is_same>::value && + std::is_same::value) returnValue += static_cast( __spv::MatrixOperands::MatrixABFloat8ComponentsINTEL) + static_cast( __spv::MatrixOperands::MatrixBBFloat8ComponentsINTEL); - if constexpr (std::is_same< - Ta, sycl::ext::oneapi::experimental::fp8_e5m2>::value && - std::is_same< - Tb, sycl::ext::oneapi::experimental::fp8_e4m3>::value && - std::is_same::value) + if constexpr ( + std::is_same>::value && + std::is_same>::value && + std::is_same::value) returnValue += static_cast( __spv::MatrixOperands::MatrixABFloat8ComponentsINTEL) + static_cast( __spv::MatrixOperands::MatrixBHFloat8ComponentsINTEL); - if constexpr (std::is_same< - Ta, sycl::ext::oneapi::experimental::fp8_e4m3>::value && - std::is_same< - Tb, sycl::ext::oneapi::experimental::fp8_e5m2>::value && - std::is_same::value) + if constexpr ( + std::is_same>::value && + std::is_same>::value && + std::is_same::value) returnValue += static_cast( __spv::MatrixOperands::MatrixAHFloat8ComponentsINTEL) + static_cast( __spv::MatrixOperands::MatrixBBFloat8ComponentsINTEL); - if constexpr (std::is_same< - Ta, sycl::ext::oneapi::experimental::fp8_e4m3>::value && - std::is_same< - Tb, sycl::ext::oneapi::experimental::fp8_e4m3>::value && - std::is_same::value) + if constexpr ( + std::is_same>::value && + std::is_same>::value && + std::is_same::value) returnValue += static_cast( __spv::MatrixOperands::MatrixAHFloat8ComponentsINTEL) + static_cast( diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp index 22642e821719e..c23885b34abee 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp @@ -589,13 +589,11 @@ joint_matrix_convert(Group, dst.spvm = __spirv_ConvertFP4E2M1ToHF16INTEL(src.spvm); else if constexpr (std::is_same::value) dst.spvm = __spirv_ConvertFP4E2M1ToBF16INTEL(src.spvm); - else if constexpr (std::is_same< - To, - sycl::ext::oneapi::experimental::fp8_e4m3>::value) + else if constexpr (std::is_same>::value) dst.spvm = __spirv_ConvertFP4E2M1ToHF8INTEL(src.spvm); - else if constexpr (std::is_same< - To, - sycl::ext::oneapi::experimental::fp8_e5m2>::value) + else if constexpr (std::is_same>::value) dst.spvm = __spirv_ConvertFP4E2M1ToBF8INTEL(src.spvm); } // FP4E2M1 down conversion @@ -623,9 +621,9 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_bmad( sycl::ext::oneapi::experimental::matrix::layout::dynamic> &D, const joint_matrix &A, const joint_matrix &B, - const joint_matrix, use::scale, M, K / 32, LayoutAs> &Ascale, - const joint_matrix, use::scale, K / 32, N, LayoutBs> &Bscale, const joint_matrix diff --git a/sycl/include/sycl/ext/oneapi/matrix/query-types.hpp b/sycl/include/sycl/ext/oneapi/matrix/query-types.hpp index d366548dd3c49..56db06e97d7cc 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/query-types.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/query-types.hpp @@ -88,14 +88,14 @@ constexpr const char *convertTypeToMatrixTypeString< return "matrix_type::tf32"; } template <> -constexpr const char * -convertTypeToMatrixTypeString() { +constexpr const char *convertTypeToMatrixTypeString< + sycl::ext::oneapi::experimental::fp8_e5m2_x<1>>() { return "matrix_type::fp8_e5m2"; } template <> -constexpr const char * -convertTypeToMatrixTypeString() { +constexpr const char *convertTypeToMatrixTypeString< + sycl::ext::oneapi::experimental::fp8_e4m3_x<1>>() { return "matrix_type::fp8_e4m3"; } diff --git a/sycl/test-e2e/Matrix/Inputs/common.hpp b/sycl/test-e2e/Matrix/Inputs/common.hpp index b870a4d591b4d..d0744e73ea491 100644 --- a/sycl/test-e2e/Matrix/Inputs/common.hpp +++ b/sycl/test-e2e/Matrix/Inputs/common.hpp @@ -11,6 +11,10 @@ #include #include #include +// The matrix headers only forward declare the fp4/fp8 types, so the tests that +// name them have to include the corresponding extension headers themselves. +#include +#include #include #include