Merge Pull Request #230 from abagusetty/EKAT/feature/sycl

Automatically Merged using E3SM Pull Request AutoTester
PR Title: Support for SYCL
PR Author: abagusetty

Author: E3SM-Bot

CMakeLists.txt

@@ -95,14 +95,18 @@ option (EKAT_DISABLE_TPL_WARNINGS "Whether we should suppress warnings when comp
 set (EKAT_TPL_LIBRARIES_INTERNAL)
 
 # A good chunk of EKAT is a utility layer over kokkos
-BuildKokkos()
-list (APPEND EKAT_TPL_LIBRARIES_INTERNAL ${Kokkos_LIBRARIES})
+# Note: for SYCL/HIP backend, bleeding edge of kokkos/develop might be needed to test some functionality
+# Kokkos_ROOT can point to custom installation directory
+if (Kokkos_ROOT)
+  find_package(Kokkos REQUIRED)
+  set(Kokkos_LIBRARIES Kokkos::kokkos)
+else()
+  BuildKokkos()
+endif()
+  list (APPEND EKAT_TPL_LIBRARIES_INTERNAL ${Kokkos_LIBRARIES})
 
 set (EKAT_ENABLE_GPU False)
-if (Kokkos_ENABLE_CUDA)
-  set (EKAT_ENABLE_GPU True)
-endif ()
-if (Kokkos_ENABLE_HIP)
+if (Kokkos_ENABLE_CUDA OR Kokkos_ENABLE_HIP OR Kokkos_ENABLE_SYCL)
   set (EKAT_ENABLE_GPU True)
 endif ()
 

cmake/EkatSetCompilerFlags.cmake

@@ -131,7 +131,7 @@ macro (SetWarningFlags)
     set (FFLAGS -Wall)
     set (CXXFLAGS -Wall)
   endif()
-    
+
   SetFlags(FFLAGS ${FFLAGS} CFLAGS -Wall CXXFLAGS ${CXXFLAGS})
 endmacro()
 
@@ -217,6 +217,8 @@ macro (SetReleaseFlags)
         string (APPEND CMAKE_CXX_FLAGS_RELEASE " -fopenmp-simd")
       elseif (CMAKE_CXX_COMPILER_ID STREQUAL "Intel")
         string (APPEND CMAKE_CXX_FLAGS_RELEASE " -qopenmp-simd")
+      elseif (CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM")
+        string (APPEND CMAKE_CXX_FLAGS_RELEASE " -fsycl -fsycl-unnamed-lambda -sycl-std=2020 -qopenmp-simd -Wsycl-strict -fsycl-device-code-split=per_kernel")
       endif()
     endif ()
   endif ()
@@ -275,6 +277,8 @@ macro (SetDebugFlags)
         string (APPEND CMAKE_CXX_FLAGS_DEBUG " -fopenmp-simd")
       elseif (CMAKE_CXX_COMPILER_ID STREQUAL "Intel")
         string (APPEND CMAKE_CXX_FLAGS_DEBUG " -qopenmp-simd")
+      elseif (CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM")
+        string (APPEND CMAKE_CXX_FLAGS_DEBUG " -qopenmp-simd -fsycl -fsycl-unnamed-lambda -sycl-std=2020 -Wsycl-strict -fno-sycl-dead-args-optimization")
       endif()
     endif ()
   endif ()

cmake/machine-files/jlse-gen9.cmake

@@ -0,0 +1,3 @@
+# Load gen9 arch with SYCL backend for kokkos
+set(CMAKE_CXX_STANDARD 17)
+include (${CMAKE_CURRENT_LIST_DIR}/kokkos/intel-gen9.cmake)

cmake/machine-files/jlse-xehp.cmake

@@ -0,0 +1,2 @@
+# Load XEHP arch with SYCL backend for kokkos
+include (${CMAKE_CURRENT_LIST_DIR}/kokkos/intel-xehp.cmake)

cmake/machine-files/kokkos/intel-gen9.cmake

@@ -0,0 +1,5 @@
+include (${CMAKE_CURRENT_LIST_DIR}/generic.cmake)
+
+option(Kokkos_ARCH_INTEL_GEN9 "" ON)
+set(Kokkos_ENABLE_SYCL TRUE CACHE BOOL "")
+set(Kokkos_ENABLE_DEPRECATED_CODE_3 FALSE CACHE BOOL "")

cmake/machine-files/kokkos/intel-xehp.cmake

@@ -0,0 +1,5 @@
+include (${CMAKE_CURRENT_LIST_DIR}/generic.cmake)
+
+option(Kokkos_ARCH_INTEL_XEHP "" ON)
+set(Kokkos_ENABLE_SYCL TRUE CACHE BOOL "")
+set(Kokkos_ENABLE_DEPRECATED_CODE_3 FALSE CACHE BOOL "")

src/ekat/ekat.hpp

@@ -25,8 +25,10 @@ static constexpr bool ekatBFB = false;
 typedef Kokkos::Cuda EkatGpuSpace;
 # elif defined KOKKOS_ENABLE_HIP
 typedef Kokkos::Experimental::HIP EkatGpuSpace;
+# elif defined KOKKOS_ENABLE_SYCL
+typedef Kokkos::Experimental::SYCL EkatGpuSpace;
 # else
-error "EKAT does not recognize a GPU space other than Cuda and HIP."
+error "EKAT does not recognize a GPU space other than Cuda, HIP and SYCL".
 # endif
 #endif
 

src/ekat/ekat_assert.hpp

@@ -27,13 +27,25 @@
     }                                                                   \
   } while(0)
 
-#define IMPL_KERNEL_THROW(condition, msg)                             \
-  do {                                                                \
-    if ( ! (condition) ) {                                            \
-      printf("KERNEL CHECK FAILED:\n   %s\n   %s\n",#condition,msg);  \
-      Kokkos::abort("");                                              \
-    }                                                                 \
+// SYCL cannot printf like the other backends quite yet
+#ifdef __SYCL_DEVICE_ONLY__
+#define IMPL_KERNEL_THROW(condition, msg)				\
+  do {									\
+    if ( ! (condition) ) {						\
+      const __attribute__((opencl_constant)) char format[] = "KERNEL CHECK FAILED:\n   %s %s\n"; \
+      sycl::ext::oneapi::experimental::printf(format,#condition,msg);	\
+      Kokkos::abort("");						\
+    }									\
   } while (0)
+#else
+#define IMPL_KERNEL_THROW(condition, msg)				\
+  do {									\
+    if ( ! (condition) ) {						\
+      printf("KERNEL CHECK FAILED:\n   %s\n   %s\n",#condition,msg);	\
+      Kokkos::abort("");						\
+    }									\
+  } while (0)
+#endif
 
 #ifndef NDEBUG
 #define EKAT_ASSERT(condition)                      IMPL_THROW(condition, "",  std::logic_error)

src/ekat/ekat_config.h.in

@@ -22,7 +22,7 @@
 // Decide whether ekat defaults to BFB behavior when possible/appropriate
 #cmakedefine EKAT_FPE
 
-// A GPU space has been enabled in Kokkos, e.g., CUDA or HIP.
+// A GPU space has been enabled in Kokkos, e.g., CUDA or HIP OR SYCL.
 #cmakedefine EKAT_ENABLE_GPU
 
 #endif // EKAT_CONFIG_H

src/ekat/ekat_session.cpp

@@ -32,6 +32,18 @@ void initialize_kokkos () {
 # elif defined KOKKOS_ENABLE_HIP
   const auto ret = hipGetDeviceCount(&nd);
   const bool ok = ret == hipSuccess;
+# elif defined KOKKOS_ENABLE_SYCL
+  nd = 0;
+  auto gpu_devs = sycl::device::get_devices(sycl::info::device_type::gpu);
+  for (auto &dev : gpu_devs) {
+    if (dev.get_info<sycl::info::device::partition_max_sub_devices>() > 0) {
+      auto subDevs = dev.create_sub_devices<sycl::info::partition_property::partition_by_affinity_domain>(sycl::info::partition_affinity_domain::numa);
+      nd += subDevs.size();
+    } else {
+      nd++;
+    }
+  }
+  const bool ok = true;
 # else
   error "No valid GPU space, yet EKAT_ENABLE_GPU is defined."
 # endif