Resolve conflicts

Author: LiyangLingIntel

include/cutlass/gemm/collective/xe_array_mma_mixed_input.hpp

@@ -173,23 +173,23 @@ struct CollectiveMma<
 
   using MmaAtomShape = typename TiledMma::AtomShape_MNK;
 
-  static constexpr auto BLK_M = get<0>(WorkgroupTileShape{});
-  static constexpr auto BLK_N = get<1>(WorkgroupTileShape{});
-  static constexpr auto BLK_K = get<2>(WorkgroupTileShape{});
+  static constexpr int BLK_M = get<0>(WorkgroupTileShape{});
+  static constexpr int BLK_N = get<1>(WorkgroupTileShape{});
+  static constexpr int BLK_K = get<2>(WorkgroupTileShape{});
 
-  static constexpr auto ATOM_M = get<1>(typename TiledMma::ThrLayoutVMNK{}.shape());
-  static constexpr auto ATOM_N = get<2>(typename TiledMma::ThrLayoutVMNK{}.shape());
-  static constexpr auto ATOM_K = get<3>(typename TiledMma::ThrLayoutVMNK{}.shape());
+  static constexpr int ATOM_M = get<1>(typename TiledMma::ThrLayoutVMNK{}.shape());
+  static constexpr int ATOM_N = get<2>(typename TiledMma::ThrLayoutVMNK{}.shape());
+  static constexpr int ATOM_K = get<3>(typename TiledMma::ThrLayoutVMNK{}.shape());
 
-  static constexpr auto SG_M = ceil_div(BLK_M, ATOM_M);
-  static constexpr auto SG_N = ceil_div(BLK_N, ATOM_N);
-  static constexpr auto SG_K = ceil_div(BLK_K, ATOM_K);
+  static constexpr int SG_M = ceil_div(BLK_M, ATOM_M);
+  static constexpr int SG_N = ceil_div(BLK_N, ATOM_N);
+  static constexpr int SG_K = ceil_div(BLK_K, ATOM_K);
   using SubgroupTileShape = Shape<decltype(SG_M), decltype(SG_N), decltype(SG_K)>;
 
   using GmemTiledCopyScale = typename scale_zero_copy_traits<NonVoidElementScale, SG_N>::type;
   using GmemTiledCopyZero = typename scale_zero_copy_traits<NonVoidElementZero, SG_N, InternalNonVoidStrideZero>::type;
 
-  static constexpr auto Num_SGs = ATOM_N * ATOM_M * ATOM_K;
+  static constexpr int Num_SGs = ATOM_N * ATOM_M * ATOM_K;
   static constexpr uint32_t MaxThreadsPerBlock = size(TiledMma{});
 
   using CopyThreadShape = Shape<_1, Int<SubgroupSize>>;

test/unit/cute/turing/movm.cu

@@ -41,8 +41,8 @@
 using namespace cute;
 
 #ifdef CUTLASS_ENABLE_SYCL
-namespace sc = syclcompat;
-namespace sc_exp = syclcompat::experimental;
+namespace sc = cutlasscompat;
+namespace sc_exp = cutlasscompat::experimental;
 namespace sycl_ext = sycl::ext::oneapi::experimental;
 #endif
 

tools/util/include/cutlass/util/initialize_block.hpp

@@ -34,7 +34,7 @@
 
 #pragma once
 #ifdef CUTLASS_ENABLE_SYCL
-#include <syclcompat.hpp> 
+#include <cutlasscompat.hpp> 
 #else
 #include <cuda.h>
 #endif
@@ -111,7 +111,7 @@ bool initialize_block(Element* block, std::size_t size, uint64_t seed, Args_t&&.
     }
   }
 
-  syclcompat::wait();
+  cutlasscompat::wait();
   return true;
 }
 
@@ -189,7 +189,7 @@ void initialize_mixed_dtype_block(cutlass::DeviceAllocation<T1>& block_device,
     }
   }
 
-  syclcompat::wait();
+  cutlasscompat::wait();
 }
 
 #undef CUDA_CHECK

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

@@ -189,8 +189,8 @@ void BlockFillRandomUniformCopyFromHost(
     for (size_t i = 0; i < capacity; ++i) {
       buff[i] = (Element)(dis(gen));
     }
-    syclcompat::memcpy<Element>(ptr, buff.data(), capacity);
-    syclcompat::wait();
+    cutlasscompat::memcpy<Element>(ptr, buff.data(), capacity);
+    cutlasscompat::wait();
   } else {
     assert(false && "Not supported dtype");
   }

tools/util/include/cutlasscompat/launch.hpp

@@ -33,8 +33,6 @@
 #include <cutlasscompat/dims.hpp>
 #include <cutlasscompat/launch_policy.hpp>
 
-#include <cute/container/type_list.hpp>
-
 namespace cutlasscompat {
 
 namespace detail {
@@ -154,7 +152,6 @@ sycl::event launch(LaunchPolicy launch_policy, sycl::queue q, Args... args) {
 template <auto F, class N=sycl::detail::auto_name, typename LaunchPolicy, typename... Args>
 sycl::event launch(LaunchPolicy launch_policy, sycl::queue q, Args... args) {
   static_assert(detail::is_launch_policy_v<LaunchPolicy>);
-  //using FN = std::conditional_t<std::is_same_v<N, sycl::detail::auto_name>, sycl::detail::auto_name, cute::type_list<decltype(F), N>>;
   return detail::launch<F, N>(launch_policy, q, args...);
 }