add 16bits x 8bits grouped gemm UTs (#495)

add 16bits x 8 bits grouped gemm UTs where test the direct conversion
from 8bits(float_e5m2) to 16bits(half_t)

Author: tdeng5

include/cutlass/gemm/collective/builders/xe_mma_builder.inl

@@ -161,8 +161,8 @@ struct CollectiveBuilder<
   KernelScheduleType,
   cute::enable_if_t<
     cute::is_any_of_v<KernelScheduleType, KernelScheduleAuto, KernelXe, KernelXeCooperative, KernelXePtrArrayCooperative> &&
-    cute::is_any_of_v<ElementA, bfloat16_t, half_t, cute::int8_t> &&
-    cute::is_any_of_v<ElementB, bfloat16_t, half_t, cute::int8_t, cute::uint4_t>
+    cute::is_any_of_v<ElementA, bfloat16_t, half_t, cute::float_e5m2_t, cute::float_e4m3_t, cute::int8_t> &&
+    cute::is_any_of_v<ElementB, bfloat16_t, half_t, cute::float_e5m2_t, cute::float_e4m3_t, cute::int8_t, cute::uint4_t>
   >
     >{
 

test/unit/gemm/device/CMakeLists.txt

@@ -79,6 +79,7 @@ if(CUTLASS_ENABLE_SYCL)
       xe_gemm_fp16_s8_fp32_tensor_op_fp32_group_gemm.cpp
       xe_gemm_bf16_u4_fp32_tensor_op_fp32_group_gemm.cpp
       xe_gemm_fp16_u4_fp32_tensor_op_fp32_group_gemm.cpp
+      xe_gemm_fp16_fp8_fp32_tensor_op_fp32_group_gemm.cpp
     )
 
     add_custom_target(

test/unit/gemm/device/default_gemm_group_configuration.hpp

@@ -71,8 +71,8 @@ struct DefaultGemmGroupConfiguration<
     ElementOutput>
 {
 
-  static_assert(cute::is_any_of_v<ElementA, bfloat16_t, half_t, int8_t>, "ElementA needs to be of 16 or 8 bit type");
-  static_assert(cute::is_any_of_v<ElementB, bfloat16_t, half_t, int8_t, uint4_t>, "ElementB needs to be of 16, 8 or 4 bit type");
+  static_assert(cute::is_any_of_v<ElementA, bfloat16_t, half_t, int8_t, float_e5m2_t, float_e4m3_t>, "ElementA needs to be of 16 or 8 bit type");
+  static_assert(cute::is_any_of_v<ElementB, bfloat16_t, half_t, int8_t, float_e5m2_t, float_e4m3_t, uint4_t>, "ElementB needs to be of 16, 8 or 4 bit type");
   using TileShape = Shape<_256, _256, _32>;
 
   using CollectiveMainloop = typename gemm::collective::CollectiveBuilder<

test/unit/gemm/device/xe_gemm_fp16_fp8_fp32_tensor_op_fp32_group_gemm.cpp

@@ -0,0 +1,101 @@
+/***************************************************************************************************
+ * Copyright (c) 2025 - 2025 Codeplay Software Ltd. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ * list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+ * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+ * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ **************************************************************************************************/
+
+/*! \file
+    \brief Tests for Xe Group fp16_fp8_fp32
+*/
+
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/kernel/gemm_universal.hpp"
+#include "cutlass/gemm/group_array_problem_shape.hpp"
+
+#include "default_gemm_configuration.hpp"
+#include "default_gemm_group_configuration.hpp"
+#include "gemm_testbed_3x_ptr_array.hpp"
+
+namespace cutlass {
+namespace {
+template <typename LayoutA, typename LayoutB>
+struct XE_Device_Gemm_fp16_fp8_f32_tensor_op_f32_group_gemm {
+  using ProblemShape = gemm::GroupProblemShape<cute::Shape<int,int,int>>; // <M,N,K> per group
+  using ElementA = cute::half_t;
+  using ElementB = cute::float_e5m2_t;
+  using ElementC = float;
+  using ElementAccumulator = float;
+  using LayoutC = layout::RowMajor;
+
+  using Config = gemm::device::DefaultGemmGroupConfiguration<
+    arch::OpClassTensorOp, arch::IntelXe,
+    ElementA, LayoutA, ElementB, LayoutB, ElementC, LayoutC, ElementAccumulator>;
+
+  using GemmKernel = gemm::kernel::GemmUniversal<
+      ProblemShape,
+      typename Config::CollectiveMainloop,
+      typename Config::CollectiveEpilogue,
+      gemm::GroupScheduler
+  >;
+
+  using Gemm = gemm::device::GemmUniversalAdapter<GemmKernel>;
+};
+
+TEST(XE_Device_Gemm_fp16t_fp8t_f32t_tensor_op_f32_group_gemm, 256x256x32) {
+  using LayoutA = layout::RowMajor;
+  using LayoutB = layout::RowMajor;
+  using Gemm = XE_Device_Gemm_fp16_fp8_f32_tensor_op_f32_group_gemm<LayoutA, LayoutB>::Gemm;
+  EXPECT_TRUE(test::gemm::device::TestAll<Gemm>(1.0, 1.0));
+  EXPECT_TRUE(test::gemm::device::TestAll<Gemm>(1.0, 0.0));
+}
+
+TEST(XE_Device_Gemm_fp16n_fp8t_f32t_tensor_op_f32_group_gemm, 256x256x32) {
+  using LayoutA = layout::ColumnMajor;
+  using LayoutB = layout::RowMajor;
+  using Gemm = XE_Device_Gemm_fp16_fp8_f32_tensor_op_f32_group_gemm<LayoutA, LayoutB>::Gemm;
+  EXPECT_TRUE(test::gemm::device::TestAll<Gemm>(1.0, 1.0));
+  EXPECT_TRUE(test::gemm::device::TestAll<Gemm>(1.0, 0.0));
+}
+
+TEST(XE_Device_Gemm_fp16t_fp8n_f32t_tensor_op_f32_group_gemm, 256x256x32) {
+  using LayoutA = layout::RowMajor;
+  using LayoutB = layout::ColumnMajor;
+  using Gemm = XE_Device_Gemm_fp16_fp8_f32_tensor_op_f32_group_gemm<LayoutA, LayoutB>::Gemm;
+  EXPECT_TRUE(test::gemm::device::TestAll<Gemm>(1.0, 1.0));
+  EXPECT_TRUE(test::gemm::device::TestAll<Gemm>(1.0, 0.0));
+}
+
+TEST(XE_Device_Gemm_fp16n_fp8n_f32t_tensor_op_f32_group_gemm, 256x256x32) {
+  using LayoutA = layout::ColumnMajor;
+  using LayoutB = layout::ColumnMajor;
+  using Gemm = XE_Device_Gemm_fp16_fp8_f32_tensor_op_f32_group_gemm<LayoutA, LayoutB>::Gemm;
+  EXPECT_TRUE(test::gemm::device::TestAll<Gemm>(1.0, 1.0));
+  EXPECT_TRUE(test::gemm::device::TestAll<Gemm>(1.0, 0.0));
+}
+}
+} // namespace cutlass