resolve conflict

build.sh

@@ -2,8 +2,8 @@ script_dir=$( cd -- "$( dirname -- "${BASH_SOURCE[0]}" )" &> /dev/null && pwd )
 cp ${script_dir}/tools/clang-format/clang-format.hook ${script_dir}/.git/hooks/pre-commit
 chmod +x ${script_dir}/.git/hooks/pre-commit
 
-# https://github.com/intel/llvm/releases/tag/nightly-2024-05-16
-sycl_compiler_path=/opt/cutlass/compiler/0516/
+# https://github.com/intel/llvm/releases/tag/nightly-2024-07-03
+sycl_compiler_path=/opt/cutlass/compiler/0703/
 
 # https://ubit-gfx.intel.com/build/19168301/artifacts
 gpu_driver_path=/opt/cutlass/gpu_driver/gfx-driver-ci-comp_igc-25012/extract/
@@ -16,10 +16,10 @@ output=intel_gpu_pvc
 unset epilogue
 
 # epilogue relu
-#epilogue+=" -DEPILOGUE_RELU "
+# epilogue+=" -DEPILOGUE_RELU "
 
 # epilogue softmax
-#epilogue+=" -DEPILOGUE_SOFTMAX "
+# epilogue+=" -DEPILOGUE_SOFTMAX "
 
 export ZE_AFFINITY_MASK=0
 export CPATH=$sycl_compiler_path:$sycl_compiler_path/include/:$sycl_compiler_path/include/sycl/

include/cutlass/epilogue/collective/default_epilogue.hpp

@@ -147,41 +147,6 @@ class DefaultEpilogue {
     return epilogue_op.is_source_needed();
   }
 
-#ifdef EPILOGUE_RELU
-  template<
-    class ProblemShapeMNKL,
-    class BlockShapeMNK,
-    class BlockCoordMNKL,
-    class FrgEngine, class FrgLayout>
-  CUTLASS_HOST_DEVICE void
-  operator()(
-    ProblemShapeMNKL problem_shape_mnkl,
-    BlockShapeMNK blk_shape_MNK,
-    BlockCoordMNKL blk_coord_mnkl,
-    cute::Tensor<FrgEngine, FrgLayout> & accumulators){
-    auto M = get<0>(problem_shape_mnkl);
-    auto N = get<1>(problem_shape_mnkl);
-    auto L = get<3>(problem_shape_mnkl);
-
-    auto [m_coord, n_coord, k_coord, l_coord] = blk_coord_mnkl;
-    if (epilogue_op.is_source_needed()) {
-      auto source = make_fragment_like(accumulators);
-      auto gmem_tiled_copy_c =
-        make_xe_2d_copy<XE_2D_U32x8x16x1x1_LD_N>(make_tensor(
-            params.ptr_C, make_shape(M, N, L), params.dC));
-
-      Tensor tCi = gmem_tiled_copy_c.get_pvc_tensor(
-          make_coord(m_coord, n_coord, l_coord),
-          make_shape(size<1>(accumulators), size<2>(accumulators), L),
-          make_stride(size<0>(blk_shape_MNK), size<1>(blk_shape_MNK)));
-      copy(gmem_tiled_copy_c, tCi(_, _, _, l_coord), source);  
-      epilogue_op(accumulators, source);
-    } else {
-      epilogue_op(accumulators);
-    } 
-  }
-#endif
-
   template<
     class ProblemShapeMNKL,
     class BlockShapeMNK,

include/cutlass/gemm/kernel/intel_pvc_gemm.hpp

@@ -197,8 +197,6 @@ class GemmUniversal<
   CUTLASS_DEVICE
   void operator()(Params const& params, char* smem_buf) {
 
-    SharedStorage& shared_storage = *reinterpret_cast<SharedStorage*>(smem_buf);
-
     // Preconditions
     CUTE_STATIC_ASSERT(is_static<TileShape>::value);