Merge branch 'main' into liyang/import_syclcompat

Author: leonling-ll

benchmarks/gemm/benchmark_runner.hpp

@@ -176,12 +176,12 @@ struct BenchmarkRunnerGemm {
 
   using CollectiveMainloop = typename Gemm::GemmKernel::CollectiveMainloop;
   using DispatchPolicy = typename CollectiveMainloop::DispatchPolicy;
-  using ElementMma = CollectiveMainloop::TiledMma::ValTypeA;
+  using ElementMma = typename CollectiveMainloop::TiledMma::ValTypeA;
 
-  using ElementScale = ScaleType<CollectiveMainloop>::type;
-  using ElementZero = ZeroType<CollectiveMainloop>::type;
-  using StrideS = ScaleStride<CollectiveMainloop>::type;
-  using StrideZ = ZeroStride<CollectiveMainloop>::type;
+  using ElementScale = typename ScaleType<CollectiveMainloop>::type;
+  using ElementZero = typename ZeroType<CollectiveMainloop>::type;
+  using StrideS = typename ScaleStride<CollectiveMainloop>::type;
+  using StrideZ = typename ZeroStride<CollectiveMainloop>::type;
 
   using CollectiveEpilogue = typename Gemm::CollectiveEpilogue;
   using ElementC = typename Gemm::ElementC;
@@ -454,7 +454,10 @@ struct BenchmarkRunnerGemm {
   }
 
   bool verify(const ProblemShapeType& problem_size, ElementCompute alpha, ElementCompute beta) {
-    auto [M, N, K, L] = problem_size;
+    auto& M = cute::get<0>(problem_size);
+    auto& N = cute::get<1>(problem_size);
+    auto& K = cute::get<2>(problem_size);
+    auto& L = cute::get<3>(problem_size);
 
     TensorRef ref_C(block_C[0].get(), LayoutC::packed({M, N}));
     TensorRef ref_D(block_ref_D.get(), LayoutD::packed({M, N}));

cmake/FindDPCPP.cmake

@@ -40,6 +40,7 @@ add_library(DPCPP::DPCPP INTERFACE IMPORTED)
 
 set(DPCPP_FLAGS "-fsycl;")
 set(DPCPP_COMPILE_ONLY_FLAGS "")
+set(DPCPP_LINK_ONLY_FLAGS "")
 
 if(NOT "${DPCPP_SYCL_TARGET}" STREQUAL "")
   list(APPEND DPCPP_FLAGS "-fsycl-targets=${DPCPP_SYCL_TARGET};")
@@ -63,10 +64,10 @@ if("${DPCPP_SYCL_TARGET}" STREQUAL "intel_gpu_pvc" OR
    "${DPCPP_SYCL_TARGET}" STREQUAL "spir64" OR
    "${DPCPP_SYCL_TARGET}" STREQUAL "intel_gpu_bmg_g21")
   if ((CMAKE_CXX_COMPILER_ID MATCHES "IntelLLVM" AND CMAKE_CXX_COMPILER_VERSION VERSION_LESS 2025.2) OR CUTLASS_SYCL_BUILTIN_ENABLE)
-    list(APPEND DPCPP_FLAGS "-Xspirv-translator;-spirv-ext=+SPV_INTEL_split_barrier")
+    list(APPEND DPCPP_LINK_ONLY_FLAGS "-Xspirv-translator;-spirv-ext=+SPV_INTEL_split_barrier")
   else()
-    list(APPEND DPCPP_FLAGS "-Xspirv-translator;-spirv-ext=+SPV_INTEL_split_barrier,+SPV_INTEL_2d_block_io,+SPV_INTEL_subgroup_matrix_multiply_accumulate")
- endif()
+    list(APPEND DPCPP_LINK_ONLY_FLAGS "-Xspirv-translator;-spirv-ext=+SPV_INTEL_split_barrier,+SPV_INTEL_2d_block_io,+SPV_INTEL_subgroup_matrix_multiply_accumulate")
+  endif()
   if(DPCPP_DISABLE_ITT_FOR_CUTLASS)
     list(APPEND DPCPP_FLAGS "-fno-sycl-instrument-device-code")
   endif()
@@ -76,14 +77,16 @@ endif()
 if(UNIX)
   set_target_properties(DPCPP::DPCPP PROPERTIES
     INTERFACE_COMPILE_OPTIONS "${DPCPP_FLAGS};${DPCPP_COMPILE_ONLY_FLAGS}"
-    INTERFACE_LINK_OPTIONS "${DPCPP_FLAGS}"
+    INTERFACE_LINK_OPTIONS "${DPCPP_FLAGS};${DPCPP_LINK_ONLY_FLAGS}"
     INTERFACE_LINK_LIBRARIES ${DPCPP_LIB_DIR}
     INTERFACE_INCLUDE_DIRECTORIES "${DPCPP_BIN_DIR}/../include/sycl;${DPCPP_BIN_DIR}/../include")
   message(STATUS "DPCPP INCLUDE DIR: ${DPCPP_BIN_DIR}/../include/sycl;${DPCPP_BIN_DIR}/../include")
-  message(STATUS "Using DPCPP flags: ${DPCPP_FLAGS};${DPCPP_COMPILE_ONLY_FLAGS}")
+  message(STATUS "Using DPCPP compile flags: ${DPCPP_FLAGS};${DPCPP_COMPILE_ONLY_FLAGS}")
+  message(STATUS "Using DPCPP link flags: ${DPCPP_FLAGS};${DPCPP_LINK_ONLY_FLAGS}")
 else()
   set_target_properties(DPCPP::DPCPP PROPERTIES
     INTERFACE_COMPILE_OPTIONS "${DPCPP_FLAGS};${DPCPP_COMPILE_ONLY_FLAGS}"
+    INTERFACE_LINK_OPTIONS "${DPCPP_FLAGS};${DPCPP_LINK_ONLY_FLAGS}"
     INTERFACE_LINK_LIBRARIES ${DPCPP_LIB_DIR}
     INTERFACE_INCLUDE_DIRECTORIES "${DPCPP_BIN_DIR}/../include/sycl")
 endif()
@@ -105,7 +108,7 @@ function(add_sycl_to_target)
   )
   get_target_property(target_type ${CUTLASS_ADD_SYCL_TARGET} TYPE)
   if (NOT target_type STREQUAL "OBJECT_LIBRARY")
-    target_link_options(${CUTLASS_ADD_SYCL_TARGET} PUBLIC ${DPCPP_FLAGS})
+    target_link_options(${CUTLASS_ADD_SYCL_TARGET} PUBLIC ${DPCPP_FLAGS} ${DPCPP_LINK_ONLY_FLAGS})
   endif()
 endfunction()
 

cmake/googletest.cmake

@@ -44,6 +44,15 @@ FetchContent_Declare(
 
 FetchContent_MakeAvailable(googletest)
 
+if (CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM")
+  if (TARGET gtest)
+   # Ignore unsupported warning flags on IntelLLVM
+    target_compile_options(gtest PRIVATE -Wno-unknown-warning-option)
+    # Show -Winline warnings, but don’t let them become errors
+    target_compile_options(gtest PRIVATE -Wno-error=inline)
+  endif()
+endif()
+
 if (MSVC)
   set(gtest_force_shared_crt ON CACHE BOOL "" FORCE)
-endif()
+endif()

examples/02_bmg_gemm_mixed_dtype/02_bmg_gemm_f16_u4_f16.cpp

@@ -453,7 +453,11 @@ struct ExampleRunner {
 
   /// Initialize operands to be used in the GEMM and reference GEMM
   void initialize(Options const& options) {
-    auto [M, N, K, L] = ProblemShapeType{options.m, options.n, options.k, options.l};
+    auto problem_shape = ProblemShapeType{options.m, options.n, options.k, options.l};
+    auto& M = cute::get<0>(problem_shape);
+    auto& N = cute::get<1>(problem_shape);
+    auto& K = cute::get<2>(problem_shape);
+    auto& L = cute::get<3>(problem_shape);
 
     auto zero_elements_packed_along_k = get<0>(StrideZero{});
     const int scale_k = cute::ceil_div(options.k, options.g);

examples/03_bmg_gemm_streamk/03_bmg_gemm_streamk.cpp

@@ -86,8 +86,6 @@ using namespace cute;
 
 ///////////////////////////////////////////////////////////////////////////////////////////////////
 
-#define CUTLASS_SYCL_PROFILING_ENABLED
-
 // Command line options parsing
 struct Options {
 
@@ -303,9 +301,9 @@ struct ExampleRunner {
     if(!passed) return cutlass::Status::kErrorInternal;
 
     if (options.iterations > 0) {
-      GPU_Clock timer;
       float elapsed_time_seconds = 0.f;
       for (int i = 0; i < options.iterations; ++i) {
+        GPU_Clock timer;
         gemm_op.initialize(arguments, workspace.get());
         timer.start();
         gemm_op.run();

examples/04_bmg_grouped_gemm/04_bmg_grouped_gemm.cpp

@@ -97,7 +97,6 @@ using ElementOutput = float;          // <- data type of elements in output matr
 
 ///////////////////////////////////////////////////////////////////////////////////////////////////
 
-#define CUTLASS_SYCL_PROFILING_ENABLED
 
 // Command line options parsing
 struct Options {

include/cutlass/gemm/collective/xe_mma_mixed_input.hpp

@@ -329,7 +329,7 @@ struct CollectiveMma<
       }();
 
       if constexpr (ModeScale) {
-        return Params{tiled_copy_a, tiled_copy_b, tiled_copy_scale, {}, args.group_size};
+        return Params{tiled_copy_a, tiled_copy_b, {tiled_copy_scale}, {}, args.group_size};
       } else {
         auto ptr_Z = [&]() {
           if constexpr (sizeof_bits_v<NonVoidElementZero> < 8) {
@@ -354,7 +354,7 @@ struct CollectiveMma<
           }
         }();
 
-        return Params{tiled_copy_a, tiled_copy_b, tiled_copy_scale, tiled_copy_zero, args.group_size};
+        return Params{tiled_copy_a, tiled_copy_b, {tiled_copy_scale}, {tiled_copy_zero}, args.group_size};
       }
     }
   }

test/unit/cute/intel_xe/mma.cpp

@@ -90,6 +90,7 @@ void gemm_device(TA const *A, TB const *B, TC *C, uint32_t m, uint32_t n,
 
 #define CUTLASS_ENABLE_DEBUG_PRINTS (0)
 
+#undef LOG_THREAD
 #define LOG_THREAD (16)
 
 #if CUTLASS_ENABLE_DEBUG_PRINTS

test/unit/cute/intel_xe/utils.hpp

@@ -53,6 +53,7 @@ using namespace cutlasscompat::experimental;
 
 #define CUTLASS_ENABLE_DEBUG_PRINTS (0)
 #define LOG_GROUP (0)
+#undef LOG_THREAD
 #define LOG_THREAD (0)
 
 template <class atype, class btype, class ctype>

tools/util/include/cutlass/util/reference/device/sycl_tensor_fill.h

@@ -1,6 +1,6 @@
 /***************************************************************************************************
- * Copyright (c) 2024 - 2024 Codeplay Software Ltd. All rights reserved.
- * Copyright (C) 2025 Intel Corporation, All rights reserved.
+* Copyright (c) 2024 - 2024 Codeplay Software Ltd. All rights reserved.
+* Copyright (c) 2025 Intel Corporation, All rights reserved.
  * SPDX-License-Identifier: BSD-3-Clause
  *
  * Redistribution and use in source and binary forms, with or without
@@ -44,6 +44,7 @@
 #include "cutlass/cutlass.h"
 #include "cutlass/complex.h"
 #include "cutlass/util/reference/device/tensor_foreach.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
 #include "cutlass/tensor_view.h"
 #include "cutlass/layout/vector.h"
 
@@ -176,25 +177,18 @@ void BlockFillRandomUniformCopyFromHost(
   size_t capacity,
   uint64_t seed,                          ///< seed for RNG
   typename RealType<Element>::Type max,   ///< upper bound of distribution
-  typename RealType<Element>::Type min    ///< lower bound for distribution
+  typename RealType<Element>::Type min,   ///< lower bound for distribution
+  int bits = -1                           ///< If non-negative, specifies number of fractional bits that
+                                          ///  are not truncated to zero. Permits reducing precision of
+                                          ///  data.
   ) {
-  
-  if constexpr(std::is_same_v<Element, float> || 
-               std::is_same_v<Element, cute::bfloat16_t> || 
-               std::is_same_v<Element, cute::half_t>) {
-    std::random_device rd;
-    std::mt19937 gen(seed);
-    std::uniform_real_distribution<float> dis(min, max);
-    auto buff = std::vector<Element>(capacity);
-
-    for (size_t i = 0; i < capacity; ++i) {
-      buff[i] = (Element)(dis(gen));
-    }
-    cutlasscompat::memcpy<Element>(ptr, buff.data(), capacity);
-    cutlasscompat::wait();
-  } else {
-    assert(false && "Not supported dtype");
-  }
+
+  auto buff = std::vector<Element>(capacity);
+
+  cutlass::reference::host::BlockFillRandomUniform(buff.data(), capacity, seed, max, min, bits);
+
+  cutlasscompat::memcpy<Element>(ptr, buff.data(), capacity);
+  cutlasscompat::wait();
 }
 
 ///////////////////////////////////////////////////////////////////////////////////////////////////