diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp
index fe005cb1b0efb..8b19136688970 100644
--- a/clang/lib/CodeGen/CodeGenTypes.cpp
+++ b/clang/lib/CodeGen/CodeGenTypes.cpp
@@ -399,6 +399,11 @@ llvm::Type *CodeGenTypes::ConvertSPVCooperativeMatrixType(RecordDecl *RD) {
       CompTy = llvm::Type::getFloatTy(getLLVMContext());
     } else if (LlvmTyName == "bfloat16") {
       CompTy = llvm::Type::getInt16Ty(getLLVMContext());
+    } 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 <typename To, typename From, std::size_t R, std::size_t C,
+          __spv::MatrixUse U,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_CooperativeMatrixKHR<To, S, R, C, U> *
+    __spirv_ConvertFP4E2M1ToHF16INTEL(
+        __spv::__spirv_CooperativeMatrixKHR<From, S, R, C, U> *Object);
+template <typename To, typename From, std::size_t R, std::size_t C,
+          __spv::MatrixUse U,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_CooperativeMatrixKHR<To, S, R, C, U> *
+    __spirv_ConvertFP4E2M1ToBF16INTEL(
+        __spv::__spirv_CooperativeMatrixKHR<From, S, R, C, U> *Object);
+template <typename To, typename From, std::size_t R, std::size_t C,
+          __spv::MatrixUse U,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_CooperativeMatrixKHR<To, S, R, C, U> *
+    __spirv_ConvertFP4E2M1ToHF8INTEL(
+        __spv::__spirv_CooperativeMatrixKHR<From, S, R, C, U> *Object);
+template <typename To, typename From, std::size_t R, std::size_t C,
+          __spv::MatrixUse U,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_CooperativeMatrixKHR<To, S, R, C, U> *
+    __spirv_ConvertFP4E2M1ToBF8INTEL(
+        __spv::__spirv_CooperativeMatrixKHR<From, S, R, C, U> *Object);
+// FP4E2M1 down conversion
+template <typename To, typename From, std::size_t R, std::size_t C,
+          __spv::MatrixUse U,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_CooperativeMatrixKHR<To, S, R, C, U> *
+    __spirv_ConvertHF16ToFP4E2M1INTEL(
+        __spv::__spirv_CooperativeMatrixKHR<From, S, R, C, U> *Object);
+template <typename To, typename From, std::size_t R, std::size_t C,
+          __spv::MatrixUse U,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_CooperativeMatrixKHR<To, S, R, C, U> *
+    __spirv_ConvertBF16ToFP4E2M1INTEL(
+        __spv::__spirv_CooperativeMatrixKHR<From, S, R, C, U> *Object);
+
+template <typename TA, typename TB, typename TC, typename TAS, typename TBS,
+          std::size_t M, std::size_t K, std::size_t N, __spv::MatrixUse UA,
+          __spv::MatrixUse UB, __spv::MatrixUse UC, __spv::MatrixUse UAS,
+          __spv::MatrixUse UBS,
+          __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
+          __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
+          __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern __DPCPP_SYCL_EXTERNAL
+    __spv::__spirv_CooperativeMatrixKHR<TC, S, M, N, UC> *
+    __spirv_CooperativeMatrixMulAddScaledINTEL(
+        __spv::__spirv_CooperativeMatrixKHR<TA, S, M, K, UA> *A,
+        __spv::__spirv_CooperativeMatrixKHR<TB, S, K, N, UB> *B,
+        __spv::__spirv_CooperativeMatrixKHR<TC, S, M, N, UC> *C,
+        __spv::__spirv_CooperativeMatrixKHR<TAS, S, M, K / 32, UAS> *Ascale,
+        __spv::__spirv_CooperativeMatrixKHR<TBS, S, K / 32, N, UBS> *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 <typename T, Scope::Flag S = Scope::Flag::Subgroup, std::size_t R = 1,
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 0e2e72b41e929..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,6 +8,7 @@
 
 #pragma once
 
+#include <cstddef>                      // size_t
 #include <optional>                     // std::optional
 #include <string_view>                  // std::string_view
 #include <sycl/__spirv/spirv_types.hpp> // __spv namespace
@@ -19,9 +20,21 @@ 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
+// <sycl/ext/oneapi/experimental/float_4bit/types.hpp> and its 8-bit counterpart
+// instead would pull both extensions into the transitive closure of
+// <sycl/sycl.hpp>, which they are deliberately not part of: code that uses
+// fp4/fp8 includes those headers explicitly.
+template <size_t N> class fp4_e2m1_x;
+template <size_t N> class fp8_e4m3_x;
+template <size_t N> class fp8_e5m2_x;
+template <size_t N> class fp8_e8m0_x;
+
 namespace matrix {
 
-enum class use { a, b, accumulator };
+enum class use { a, b, accumulator, scale };
 
 enum class layout {
   row_major = 0,
@@ -49,6 +62,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 +99,11 @@ extern "C" constexpr __spv::MatrixLayout joint_matrix_layout_to_spv(
   }
 }
 
+template <typename T> struct is_fp4_e2m1 : std::false_type {};
+template <size_t N>
+struct is_fp4_e2m1<sycl::ext::oneapi::experimental::fp4_e2m1_x<N>>
+    : std::true_type {};
+
 template <typename Ta, typename Tb, typename Tc, typename Td>
 constexpr uint32_t CalculateMatrixOperand() {
   uint32_t returnValue = 0x00;
@@ -104,6 +123,44 @@ constexpr uint32_t CalculateMatrixOperand() {
   if constexpr (std::is_signed<Tb>::value)
     returnValue += static_cast<uint32_t>(
         __spv::MatrixOperands::MatrixBSignedComponentsKHR);
+  if constexpr (
+      std::is_same<Ta, sycl::ext::oneapi::experimental::fp8_e5m2_x<1>>::value &&
+      std::is_same<Tb, sycl::ext::oneapi::experimental::fp8_e5m2_x<1>>::value &&
+      std::is_same<Tc, float>::value)
+    returnValue += static_cast<uint32_t>(
+                       __spv::MatrixOperands::MatrixABFloat8ComponentsINTEL) +
+                   static_cast<uint32_t>(
+                       __spv::MatrixOperands::MatrixBBFloat8ComponentsINTEL);
+  if constexpr (
+      std::is_same<Ta, sycl::ext::oneapi::experimental::fp8_e5m2_x<1>>::value &&
+      std::is_same<Tb, sycl::ext::oneapi::experimental::fp8_e4m3_x<1>>::value &&
+      std::is_same<Tc, float>::value)
+    returnValue += static_cast<uint32_t>(
+                       __spv::MatrixOperands::MatrixABFloat8ComponentsINTEL) +
+                   static_cast<uint32_t>(
+                       __spv::MatrixOperands::MatrixBHFloat8ComponentsINTEL);
+  if constexpr (
+      std::is_same<Ta, sycl::ext::oneapi::experimental::fp8_e4m3_x<1>>::value &&
+      std::is_same<Tb, sycl::ext::oneapi::experimental::fp8_e5m2_x<1>>::value &&
+      std::is_same<Tc, float>::value)
+    returnValue += static_cast<uint32_t>(
+                       __spv::MatrixOperands::MatrixAHFloat8ComponentsINTEL) +
+                   static_cast<uint32_t>(
+                       __spv::MatrixOperands::MatrixBBFloat8ComponentsINTEL);
+  if constexpr (
+      std::is_same<Ta, sycl::ext::oneapi::experimental::fp8_e4m3_x<1>>::value &&
+      std::is_same<Tb, sycl::ext::oneapi::experimental::fp8_e4m3_x<1>>::value &&
+      std::is_same<Tc, float>::value)
+    returnValue += static_cast<uint32_t>(
+                       __spv::MatrixOperands::MatrixAHFloat8ComponentsINTEL) +
+                   static_cast<uint32_t>(
+                       __spv::MatrixOperands::MatrixBHFloat8ComponentsINTEL);
+  if constexpr (sycl::detail::is_fp4_e2m1<Ta>::value &&
+                sycl::detail::is_fp4_e2m1<Tb>::value)
+    returnValue += static_cast<uint32_t>(
+                       __spv::MatrixOperands::MatrixAFP4S1E2M1ComponentsINTEL) +
+                   static_cast<uint32_t>(
+                       __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..c23885b34abee 100644
--- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp
+++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp
@@ -576,6 +576,75 @@ joint_matrix_prefetch(Group sg, T *Ptr, size_t stride,
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
+template <typename To, typename From, typename Group, std::size_t M,
+          std::size_t N, use Use, layout Layout>
+inline __SYCL_ALWAYS_INLINE void
+joint_matrix_convert(Group,
+                     const joint_matrix<Group, From, Use, M, N, Layout> &src,
+                     joint_matrix<Group, To, Use, M, N, Layout> &dst) {
+#if defined(__SYCL_DEVICE_ONLY__)
+  // FP4E2M1 Upconversion
+  if constexpr (sycl::detail::is_fp4_e2m1<From>::value) {
+    if constexpr (std::is_same<To, sycl::half>::value)
+      dst.spvm = __spirv_ConvertFP4E2M1ToHF16INTEL<To>(src.spvm);
+    else if constexpr (std::is_same<To, sycl::ext::oneapi::bfloat16>::value)
+      dst.spvm = __spirv_ConvertFP4E2M1ToBF16INTEL<To>(src.spvm);
+    else if constexpr (std::is_same<To, sycl::ext::oneapi::experimental::
+                                            fp8_e4m3_x<1>>::value)
+      dst.spvm = __spirv_ConvertFP4E2M1ToHF8INTEL<To>(src.spvm);
+    else if constexpr (std::is_same<To, sycl::ext::oneapi::experimental::
+                                            fp8_e5m2_x<1>>::value)
+      dst.spvm = __spirv_ConvertFP4E2M1ToBF8INTEL<To>(src.spvm);
+  }
+  // FP4E2M1 down conversion
+  else if constexpr (sycl::detail::is_fp4_e2m1<To>::value) {
+    if constexpr (std::is_same<From, sycl::half>::value)
+      dst.spvm = __spirv_ConvertHF16ToFP4E2M1INTEL<To>(src.spvm);
+    else if constexpr (std::is_same<From, sycl::ext::oneapi::bfloat16>::value)
+      dst.spvm = __spirv_ConvertBF16ToFP4E2M1INTEL<To>(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 <typename Group, typename Ta, typename Tb, typename Tc, typename Td,
+          std::size_t M, std::size_t K, std::size_t N, layout LayoutA,
+          layout LayoutB, layout LayoutAs, layout LayoutBs>
+inline __SYCL_ALWAYS_INLINE void joint_matrix_bmad(
+    Group,
+    joint_matrix<Group, Td, use::accumulator, M, N,
+                 sycl::ext::oneapi::experimental::matrix::layout::dynamic> &D,
+    const joint_matrix<Group, Ta, use::a, M, K, LayoutA> &A,
+    const joint_matrix<Group, Tb, use::b, K, N, LayoutB> &B,
+    const joint_matrix<Group, sycl::ext::oneapi::experimental::fp8_e8m0_x<1>,
+                       use::scale, M, K / 32, LayoutAs> &Ascale,
+    const joint_matrix<Group, sycl::ext::oneapi::experimental::fp8_e8m0_x<1>,
+                       use::scale, K / 32, N, LayoutBs> &Bscale,
+    const joint_matrix<Group, Tc, use::accumulator, M, N,
+                       sycl::ext::oneapi::experimental::matrix::layout::dynamic>
+        &C) {
+#if defined(__SYCL_DEVICE_ONLY__)
+  constexpr uint32_t MatrixOperand =
+      sycl::detail::CalculateMatrixOperand<Ta, Tb, Tc, Td>();
+  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..56db06e97d7cc 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 <typename T> constexpr const char *convertTypeToMatrixTypeString() {
-  return "";
+  // fp4_e2m1_x<N> 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<T>::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<
+    sycl::ext::oneapi::experimental::fp8_e5m2_x<1>>() {
+  return "matrix_type::fp8_e5m2";
+}
+
+template <>
+constexpr const char *convertTypeToMatrixTypeString<
+    sycl::ext::oneapi::experimental::fp8_e4m3_x<1>>() {
+  return "matrix_type::fp8_e4m3";
+}
+
 template <> constexpr const char *convertTypeToMatrixTypeString<float>() {
   return "matrix_type::fp32";
 }
diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp
index ea45d51058753..d1588af1eb243 100644
--- a/sycl/source/detail/device_impl.hpp
+++ b/sycl/source/detail/device_impl.hpp
@@ -2206,6 +2206,33 @@ 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.
+      bool SupportsFP8AndFP4 = (architecture::intel_gpu_cri == DeviceArch);
+      if (SupportsFP8AndFP4)
+        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<sycl::exception> 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..d0744e73ea491 100644
--- a/sycl/test-e2e/Matrix/Inputs/common.hpp
+++ b/sycl/test-e2e/Matrix/Inputs/common.hpp
@@ -11,6 +11,10 @@
 #include <iostream>
 #include <random>
 #include <sycl/detail/core.hpp>
+// 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 <sycl/ext/oneapi/experimental/float_4bit/types.hpp>
+#include <sycl/ext/oneapi/experimental/float_8bit/types.hpp>
 #include <sycl/ext/oneapi/matrix/matrix.hpp>
 #include <sycl/kernel_bundle.hpp>
 
@@ -177,6 +181,159 @@ 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 <typename T>
+constexpr bool is_fp8_type_v =
+    std::is_same_v<T, syclex::fp8_e5m2> ||
+    std::is_same_v<T, syclex::fp8_e8m0> || std::is_same_v<T, syclex::fp8_e4m3>;
+
+// 4-bit types: they pack numElems elements and convert to/from
+// marray<half, numElems>.
+template <typename T, size_t numElems>
+constexpr bool is_4bit_type_v = std::is_same_v<T, syclex::fp4_e2m1_x<numElems>>;
+
+// numElems is the packing factor for the 4bits types. It can be 2 or 8
+template <typename Ts, typename Td, size_t numElems = 2>
+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<Td, sycl::half> && is_fp8_type_v<Ts>)
+           dst[i * cols + j] = (Td)src[i * cols + j];
+         else if constexpr (std::is_same_v<Ts, sycl::half> && is_fp8_type_v<Td>)
+           dst[i * cols + j] = src[i * cols + j];
+         else if constexpr (std::is_same_v<Td, sycl::half> &&
+                            is_4bit_type_v<Ts, numElems>) {
+           // src[i][j] packs numElems consecutive values, so it expands to
+           // dst[i][j * numElems .. j * numElems + numElems - 1].
+           marray<Td, numElems> mval = (marray<Td, numElems>)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<Ts, sycl::half> &&
+                              is_4bit_type_v<Td, numElems>) {
+           marray<sycl::half, numElems> 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 <typename Tprecision, typename Tin>
+Tin scalar_truncate_fraction_bits(Tin number) {
+  static_assert(std::is_same_v<Tin, sycl::half> &&
+                "Unsupported input type in scalar_truncate_fraction_bits");
+
+  union {
+    Tin as_input;
+    uint16_t as_integer;
+  } caster = {number};
+
+  if constexpr (std::is_same_v<Tprecision, syclex::fp8_e5m2>) {
+    caster.as_integer &= 0xFF00;
+  } else if constexpr (std::is_same_v<Tprecision, syclex::fp8_e4m3>) {
+    caster.as_integer &= 0xFF80;
+  } else
+    static_assert(
+        false &&
+        "Unsupported precision type in scalar_truncate_fraction_bits.");
+
+  return caster.as_input;
+}
+
+template <typename Tprecision, typename Tin>
+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<Tprecision, Tin>(mat[i * cols + j]);
+    }
+  }
+}
+
+template <typename T, typename F>
+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<T>)
+           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 <typename Ta, typename Tb, typename Tc, unsigned int VF = 1,
+          typename F = std::nullptr_t>
+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<Ta, bfloat16> &&
+                          std::is_same_v<Tc, float>)
+              acc += make_fp32(va[i]) * make_fp32(vb[i]);
+            else if constexpr (std::is_same_v<Ta, sycl::half> &&
+                               std::is_same_v<Tc, float>)
+              acc += (float)va[i] * (float)vb[i];
+            else if constexpr (std::is_same_v<Ta, float> &&
+                                   std::is_same_v<Tc, float> ||
+                               std::is_integral_v<Ta> &&
+                                   std::is_integral_v<Tc> ||
+                               (std::is_same_v<Ta, bfloat16> ||
+                                std::is_same_v<Ta, sycl::half>) ||
+                               (std::is_same_v<Ta, double> &&
+                                std::is_same_v<Tc, double>))
+              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<F, std::nullptr_t>) {
+        lambda(accb);
+      }
+      *(C + c_ind) = accb;
+    }
+  }
+}
+
 template <typename F, typename T>
 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..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
@@ -1,4 +1,5 @@
 #include <iostream>
+#include <sycl/usm.hpp>
 //==----------- element_wise_all_ops_impl.hpp  - DPC++ joint_matrix---------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
@@ -22,6 +23,21 @@ void assert_ops_ref(host_accessor<T, 2, access_mode::read> mat,
     }
 }
 
+// Overload for results held in a USM allocation instead of a buffer.
+template <typename T, size_t NUM_ROWS, size_t NUM_COLS>
+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<T, bfloat16>)
+        diff = make_fp32(mat[i * NUM_COLS + j]) - ref;
+      else
+        diff = mat[i * NUM_COLS + j] - ref;
+      assert(std::fabs(static_cast<float>(diff)) <
+             std::numeric_limits<float>::epsilon());
+    }
+}
+
 template <typename T, size_t NUM_ROWS, size_t NUM_COLS, size_t SUB_ROWS,
           size_t SUB_COLS, use Use, layout Layout, size_t VF, class kernel_name,
           typename OP>
@@ -31,7 +47,6 @@ void verify_op_ab(const T l, const T r, const float ref, OP op) {
 
   buffer<T, 2> bufMat(big_mat.get_data(),
                       range<2>(NUM_ROWS / VF, NUM_COLS * VF));
-
   queue q;
   size_t sg_size = get_sg_size<kernel_name>(q);
   q.submit([&](handler &cgh) {
@@ -65,6 +80,115 @@ void verify_op_ab(const T l, const T r, const float ref, OP op) {
       bufMat.get_host_access(read_only), ref);
 }
 
+template <typename T, size_t NUM_ROWS, size_t NUM_COLS, size_t SUB_ROWS,
+          size_t SUB_COLS, use Use, layout Layout, size_t VF, class kernel_name,
+          typename OP, std::enable_if_t<is_fp8_type_v<T>, bool> = true>
+void verify_op_ab(const sycl::half l, const sycl::half r, const float ref,
+                  OP op) {
+  queue q;
+  T *mat = malloc_shared<T>(NUM_ROWS / VF * NUM_COLS * VF, q);
+  sycl::half *matH =
+      malloc_shared<sycl::half>(NUM_ROWS / VF * NUM_COLS * VF, q);
+  size_t sg_size = get_sg_size<kernel_name>(q);
+  q.submit([&](handler &cgh) {
+     cgh.parallel_for<kernel_name>(
+         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);
+
+           auto pMat =
+               address_space_cast<sycl::access::address_space::global_space,
+                                  access::decorated::no>(mat);
+
+           sub_group sg = spmd_item.get_sub_group();
+           joint_matrix<sub_group, T, Use, SUB_ROWS, SUB_COLS, Layout> sub_mat;
+           joint_matrix_fill(sg, sub_mat, T(l));
+           joint_matrix_apply(sg, sub_mat, [=](T &x) { x = op((half)x, r); });
+           ext::intel::experimental::matrix::joint_matrix_store(
+               sg, sub_mat,
+               pMat + (sg_startx * SUB_ROWS / VF) * NUM_COLS * VF +
+                   sg_starty / sg_size * SUB_COLS * VF,
+               NUM_COLS * VF);
+         }); // parallel for
+   }).wait();
+
+  matrix_copy<T, sycl::half>(q, NUM_ROWS, NUM_COLS, mat, matH);
+
+  assert_ops_ref<sycl::half, NUM_ROWS / VF, NUM_COLS * VF>(matH, ref);
+
+  free(mat, q);
+  free(matH, q);
+}
+
+template <typename T, size_t NUM_ROWS, size_t NUM_COLS, size_t SUB_ROWS,
+          size_t SUB_COLS, use Use, layout Layout, size_t VF, class kernel_name,
+          size_t numElems = 2, typename OP,
+          std::enable_if_t<is_4bit_type_v<T, numElems>, 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<numElems> 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<T>(NUM_ROWS / VF * Cols, q);
+  sycl::half *matH =
+      malloc_shared<sycl::half>(NUM_ROWS / VF * NUM_COLS * VF, q);
+  size_t sg_size = get_sg_size<kernel_name>(q);
+  q.submit([&](handler &cgh) {
+     cgh.parallel_for<kernel_name>(
+         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);
+
+           auto pMat =
+               address_space_cast<sycl::access::address_space::global_space,
+                                  access::decorated::no>(mat);
+
+           sub_group sg = spmd_item.get_sub_group();
+           joint_matrix<sub_group, T, Use, SUB_ROWS, SUB_COLS, Layout> sub_mat;
+           marray<sycl::half, numElems> fillVal(l);
+           joint_matrix_fill(sg, sub_mat, T(fillVal));
+           joint_matrix_apply(sg, sub_mat, [=](T &x) {
+             marray<sycl::half, numElems> mval =
+                 (marray<sycl::half, numElems>)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);
+           });
+           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<T, sycl::half, numElems>(q, NUM_ROWS / VF, Cols, mat, matH);
+
+  assert_ops_ref<sycl::half, NUM_ROWS / VF, NUM_COLS * VF>(matH, ref);
+
+  free(mat, q);
+  free(matH, q);
+}
+
 template <typename T, size_t NUM_ROWS, size_t NUM_COLS, size_t SUB_ROWS,
           size_t SUB_COLS, class kernel_name, typename OP>
 void verify_op_c(const T l, const T r, const float ref, OP op) {
@@ -120,7 +244,6 @@ void test_ewops_ab() {
 
   static constexpr size_t NROWS = SROWS * 2;
   static constexpr size_t NCOLS = SCOLS * 2;
-
   verify_op_ab<T, NROWS, NCOLS, SROWS, SCOLS, Use, Layout, VF,
                ewops_ab<T, SROWS, SCOLS, Use, class ab_add>>(
       Tv(5.0), Tv(2.0), 7.0, [](auto l, auto r) { return l + r; });
@@ -252,6 +375,30 @@ 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;
+  if (is_type_supported_by_device(q, matrix_type::fp4_e2m1)) {
+    test_ewops_ab<syclex::fp4_e2m1_x<numElems>, 8, 32, use::a,
+                  layout::row_major, 1, sycl::half>();
+    test_ewops_ab<syclex::fp4_e2m1_x<numElems>, 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<syclex::fp8_e5m2, 8, 32, use::a, layout::row_major, 1,
+                  sycl::half>();
+    test_ewops_ab<syclex::fp8_e4m3, 8, 32, use::a, layout::row_major, 1,
+                  sycl::half>();
+    test_ewops_ab<syclex::fp8_e5m2, 32, 16, use::b, layout::ext_intel_packed, 4,
+                  sycl::half>();
+    test_ewops_ab<syclex::fp8_e4m3, 32, 16, use::b, layout::ext_intel_packed, 4,
+                  sycl::half>();
+  }
 
   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..015103a97728c
--- /dev/null
+++ b/sycl/test-e2e/Matrix/Inputs/joint_matrix_float4_impl.hpp
@@ -0,0 +1,212 @@
+//==---------- 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 <algorithm>
+#include <sycl/usm.hpp>
+
+constexpr size_t TM = 8;
+constexpr size_t TN = 16;
+constexpr size_t TK = 32;
+
+// numElems is the packing factor of fp4_e2m1_x<numElems>: 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 <typename TA, typename TB, typename TC, size_t M, size_t N, size_t K,
+          layout B_layout, unsigned int vnniFactor, unsigned int numElems,
+          bool convertP>
+void joint_matrix_gemm_vnni(
+    sub_group sg, size_t sg_startx, size_t sg_starty, size_t sg_size,
+    multi_ptr<TA, sycl::access::address_space::global_space,
+              access::decorated::no>
+        pA,
+    multi_ptr<std::conditional_t<convertP, sycl::half, TB>,
+              sycl::access::address_space::global_space, access::decorated::no>
+        pB,
+    multi_ptr<TC, sycl::access::address_space::global_space,
+              access::decorated::no>
+        pC) {
+  joint_matrix<sub_group, TA, use::a, TM, TK, layout::row_major> sub_a;
+  joint_matrix<sub_group, TB, use::b, TK, TN, B_layout> sub_b;
+  joint_matrix<sub_group, TC, use::accumulator, TM, TN> 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_group, sycl::half, use::b, TK, TN, B_layout> 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 <typename TA, typename TB, typename TC, unsigned int vnniFactor,
+          layout B_layout, bool convertP>
+class fp4matrix;
+
+template <typename TA, typename TB, typename TC, size_t M, size_t N, size_t K,
+          layout B_layout, unsigned int vnniFactor, unsigned int numElems,
+          bool convertP>
+void matrix_multiply(TC *C, TA *A,
+                     std::conditional_t<convertP, sycl::half, TB> *B, queue q) {
+  size_t NDRangeM = M / TM;
+  size_t NDRangeN = N / TN;
+
+  auto pA = address_space_cast<sycl::access::address_space::global_space,
+                               access::decorated::no>(A);
+  auto pB = address_space_cast<sycl::access::address_space::global_space,
+                               access::decorated::no>(B);
+  auto pC = address_space_cast<sycl::access::address_space::global_space,
+                               access::decorated::no>(C);
+  using kernel_name = fp4matrix<TA, TB, TC, vnniFactor, B_layout, convertP>;
+  size_t sg_size = get_sg_size<kernel_name>(q);
+  q.submit([&](handler &cgh) {
+     cgh.parallel_for<kernel_name>(
+         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<TA, TB, TC, M, N, K, B_layout, vnniFactor,
+                                  numElems, convertP>(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 <typename T4bit, unsigned int numElems>
+void matrix_round_trip(queue q, unsigned int rows, unsigned int packedCols,
+                       sycl::half *Mat, T4bit *Packed) {
+  matrix_copy<sycl::half, T4bit, numElems>(q, rows, packedCols, Mat, Packed);
+  matrix_copy<T4bit, sycl::half, numElems>(q, rows, packedCols, Packed, Mat);
+}
+
+template <typename TA, typename TB, typename TC, size_t M, size_t N, size_t K,
+          unsigned int vnniFactor, bool convertP, unsigned int numElems>
+void joint_matrix_verify(queue q) {
+  sycl::half *Ah = malloc_shared<sycl::half>(M * K, q);
+  sycl::half *Bh = malloc_shared<sycl::half>(K * N, q);
+  TA *A = malloc_shared<TA>(M * K / numElems, q);
+  TB *B = malloc_shared<TB>(K * N / numElems, q);
+  TC *C = malloc_shared<TC>(M * N, q);
+  TC *D = malloc_shared<TC>(M * N, q);
+
+  matrix_rand<sycl::half>(M, K, Ah, 5);
+  matrix_rand<sycl::half>(K, N, Bh, 5);
+  // Snap the reference data to values E2M1 represents exactly.
+  matrix_round_trip<TA, numElems>(q, M, K / numElems, Ah, A);
+  matrix_round_trip<TB, numElems>(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<sycl::half>(K * N, q);
+    matrix_vnni(K, N, Bh, vnniBh, vnniFactor);
+    if constexpr (convertP) {
+      matrix_multiply<TA, TB, TC, M, N, K, layout::ext_intel_packed, vnniFactor,
+                      numElems, convertP>(C, A, vnniBh, q);
+    } else {
+      TB *vnniB = malloc_shared<TB>(K * N / numElems, q);
+      matrix_copy<sycl::half, TB, numElems>(q, K, N / numElems, vnniBh, vnniB);
+      matrix_multiply<TA, TB, TC, M, N, K, layout::ext_intel_packed, vnniFactor,
+                      numElems, convertP>(C, A, vnniB, q);
+      free(vnniB, q);
+    }
+    free(vnniBh, q);
+  } else { // row major
+    if constexpr (convertP) {
+      matrix_multiply<TA, TB, TC, M, N, K, layout::row_major, vnniFactor,
+                      numElems, convertP>(C, A, Bh, q);
+    } else {
+      matrix_multiply<TA, TB, TC, M, N, K, layout::row_major, vnniFactor,
+                      numElems, convertP>(C, A, B, q);
+    }
+  }
+  matrix_multiply_ref<sycl::half, sycl::half, TC>(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 <typename TC, unsigned int numElems> 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<size_t>(TN * SCALE, 64);
+  static constexpr size_t MATRIX_K = TK * SCALE;
+
+  using fp4 = syclex::fp4_e2m1_x<numElems>;
+
+  // vnniFactor 8 fills a 32-bit dword with 4-bit elements
+  joint_matrix_verify<fp4, fp4, TC, MATRIX_M, MATRIX_N, MATRIX_K,
+                      /*vnniFactor=*/8, /*convertP=*/false, numElems>(q);
+  joint_matrix_verify<fp4, fp4, TC, MATRIX_M, MATRIX_N, MATRIX_K, 1, false,
+                      numElems>(q);
+  joint_matrix_verify<fp4, fp4, TC, MATRIX_M, MATRIX_N, MATRIX_K, 8, true,
+                      numElems>(q);
+  joint_matrix_verify<fp4, fp4, TC, MATRIX_M, MATRIX_N, MATRIX_K, 1, true,
+                      numElems>(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
+  constexpr unsigned int numElems = 2;
+  fp4_combinations<float, numElems>(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<bfloat16, numElems>(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..4a2d5eab71f79
--- /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 <algorithm>
+#include <sycl/usm.hpp>
+
+constexpr size_t TN = 16;
+constexpr size_t TK = 32;
+
+template <typename TA, typename TB, typename TC, size_t M, size_t N, size_t K,
+          size_t TM, layout B_layout, unsigned int vnniFactor>
+void joint_matrix_gemm_vnni(
+    sub_group sg, size_t sg_startx, size_t sg_starty, size_t sg_size,
+    multi_ptr<TA, sycl::access::address_space::global_space,
+              access::decorated::no>
+        pA,
+    multi_ptr<TB, sycl::access::address_space::global_space,
+              access::decorated::no>
+        pB,
+    multi_ptr<TC, sycl::access::address_space::global_space,
+              access::decorated::no>
+        pC) {
+  joint_matrix<sub_group, TA, use::a, TM, TK, layout::row_major> sub_a;
+  joint_matrix<sub_group, TB, use::b, TK, TN, B_layout> sub_b;
+  joint_matrix<sub_group, TC, use::accumulator, TM, TN> 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 <typename TA, typename TB, typename TC, size_t TM,
+          unsigned int vnniFactor, layout B_layout>
+class fp8matrix;
+
+template <typename TA, typename TB, typename TC, size_t M, size_t N, size_t K,
+          size_t TM, layout B_layout, unsigned int vnniFactor>
+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<sycl::access::address_space::global_space,
+                               access::decorated::no>(A);
+  auto pB = address_space_cast<sycl::access::address_space::global_space,
+                               access::decorated::no>(B);
+  auto pC = address_space_cast<sycl::access::address_space::global_space,
+                               access::decorated::no>(C);
+  size_t sg_size =
+      get_sg_size<fp8matrix<TA, TB, TC, TM, vnniFactor, B_layout>>(q);
+  q.submit([&](handler &cgh) {
+     cgh.parallel_for<fp8matrix<TA, TB, TC, TM, vnniFactor, B_layout>>(
+         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<TA, TB, TC, M, N, K, TM, B_layout,
+                                  vnniFactor>(sg, sg_startx, sg_starty, sg_size,
+                                              pA, pB, pC);
+         }); // parallel for
+   }).wait();
+}
+
+template <typename TA, typename TB, typename TC, size_t M, size_t N, size_t K,
+          size_t TM, unsigned int vnniFactor>
+void joint_matrix_verify(queue q) {
+  sycl::half *Ah = malloc_shared<sycl::half>(M * K, q);
+  sycl::half *Bh = malloc_shared<sycl::half>(K * N, q);
+  TA *A = malloc_shared<TA>(M * K, q);
+  TB *B = malloc_shared<TB>(K * N, q);
+  TC *C = malloc_shared<TC>(M * N, q);
+  TC *D = malloc_shared<TC>(M * N, q);
+
+  matrix_rand<sycl::half>(M, K, Ah, 5);
+  matrix_truncate_fraction_bits<TA>(M, K, Ah);
+  matrix_rand<sycl::half>(K, N, Bh, 5);
+  matrix_truncate_fraction_bits<TB>(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<TB>(K * N, q);
+    matrix_vnni<TB>(K, N, B, vnniB, vnniFactor);
+    matrix_multiply<TA, TB, TC, M, N, K, TM, layout::ext_intel_packed,
+                    vnniFactor>(C, A, vnniB, q);
+  } else {
+    matrix_multiply<TA, TB, TC, M, N, K, TM, layout::row_major, vnniFactor>(
+        C, A, B, q);
+  }
+  matrix_multiply_ref<sycl::half, sycl::half, TC>(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 <size_t TM> 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<size_t>(TN * SCALE, 64);
+  static constexpr size_t MATRIX_K = TK * SCALE;
+
+  joint_matrix_verify<syclex::fp8_e5m2, syclex::fp8_e5m2, float, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 4>(q);
+  joint_matrix_verify<syclex::fp8_e5m2, syclex::fp8_e5m2, float, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 1>(q);
+  joint_matrix_verify<syclex::fp8_e4m3, syclex::fp8_e4m3, float, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 4>(q);
+  joint_matrix_verify<syclex::fp8_e4m3, syclex::fp8_e4m3, float, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 1>(q);
+  joint_matrix_verify<syclex::fp8_e5m2, syclex::fp8_e4m3, float, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 4>(q);
+  joint_matrix_verify<syclex::fp8_e5m2, syclex::fp8_e4m3, float, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 1>(q);
+  joint_matrix_verify<syclex::fp8_e4m3, syclex::fp8_e5m2, float, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 4>(q);
+  joint_matrix_verify<syclex::fp8_e4m3, syclex::fp8_e5m2, float, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 1>(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<syclex::fp8_e5m2, syclex::fp8_e5m2, bfloat16, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 4>(q);
+  joint_matrix_verify<syclex::fp8_e5m2, syclex::fp8_e5m2, bfloat16, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 1>(q);
+  joint_matrix_verify<syclex::fp8_e4m3, syclex::fp8_e4m3, bfloat16, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 4>(q);
+  joint_matrix_verify<syclex::fp8_e4m3, syclex::fp8_e4m3, bfloat16, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 1>(q);
+  joint_matrix_verify<syclex::fp8_e5m2, syclex::fp8_e4m3, bfloat16, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 4>(q);
+  joint_matrix_verify<syclex::fp8_e5m2, syclex::fp8_e4m3, bfloat16, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 1>(q);
+  joint_matrix_verify<syclex::fp8_e4m3, syclex::fp8_e5m2, bfloat16, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 4>(q);
+  joint_matrix_verify<syclex::fp8_e4m3, syclex::fp8_e5m2, bfloat16, MATRIX_M,
+                      MATRIX_N, MATRIX_K, TM, 1>(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..0bff8c613baa6
--- /dev/null
+++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_float4.cpp
@@ -0,0 +1,22 @@
+//==----------- 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
+
+// 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
+
+// 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..cac8b48eca9bf
--- /dev/null
+++ b/sycl/test-e2e/Matrix/SG32/joint_matrix_float8.cpp
@@ -0,0 +1,19 @@
+//==----------- 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
+
+// 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/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/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"
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..7403f31b34e9d
--- /dev/null
+++ b/sycl/test-e2e/Matrix/joint_matrix_float4.cpp
@@ -0,0 +1,21 @@
+//==----------- 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
+
+// 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
+
+// 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..59824d20aa06f
--- /dev/null
+++ b/sycl/test-e2e/Matrix/joint_matrix_float8.cpp
@@ -0,0 +1,18 @@
+//==----------- 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
+
+// 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"