SYCL: Use SYCL_EXT_ONEAPI_DEVICE_GLOBAL to detect support for device …

…global variables
fnrizzi · Oct 18, 2023 · 2bc1721 · 2bc1721
1 parent 0c0cafa
commit 2bc1721
Show file tree

Hide file tree

Showing 2 changed files with 32 additions and 22 deletions.
diff --git a/cmake/kokkos_arch.cmake b/cmake/kokkos_arch.cmake
@@ -588,32 +588,39 @@ IF (KOKKOS_ENABLE_SYCL)
 ENDIF()
 
 # Check support for device_global variables
-# FIXME_SYCL Once the feature test macro SYCL_EXT_ONEAPI_DEVICE_GLOBAL is
-#            available, use that instead.
+# FIXME_SYCL Even if SYCL_EXT_ONEAPI_DEVICE_GLOBAL is defined, we still can't
+#            use device global variables with shared libraries
 IF(KOKKOS_ENABLE_SYCL AND NOT BUILD_SHARED_LIBS)
-  INCLUDE(CheckCXXSourceCompiles)
   STRING(REPLACE ";" " " CMAKE_REQUIRED_FLAGS "${KOKKOS_COMPILE_OPTIONS}")
-  CHECK_CXX_SOURCE_COMPILES("
-    #include <sycl/sycl.hpp>
-    using namespace sycl::ext::oneapi::experimental;
-    using namespace sycl;
+  INCLUDE(CheckCXXSymbolExists)
+  CHECK_CXX_SYMBOL_EXISTS(SYCL_EXT_ONEAPI_DEVICE_GLOBAL "sycl/sycl.hpp" KOKKOS_IMPL_HAVE_SYCL_EXT_ONEAPI_DEVICE_GLOBAL)
+  IF (KOKKOS_IMPL_HAVE_SYCL_EXT_ONEAPI_DEVICE_GLOBAL)
+    SET(KOKKOS_IMPL_SYCL_DEVICE_GLOBAL_SUPPORTED ON)
+    COMPILER_SPECIFIC_FLAGS(DEFAULT -DDESUL_SYCL_DEVICE_GLOBAL_SUPPORTED)
+  ELSE()
+    INCLUDE(CheckCXXSourceCompiles)
+    CHECK_CXX_SOURCE_COMPILES("
+      #include <sycl/sycl.hpp>
+      using namespace sycl::ext::oneapi::experimental;
+      using namespace sycl;
 
-    SYCL_EXTERNAL device_global<int, decltype(properties(device_image_scope))> Foo;
+      SYCL_EXTERNAL device_global<int, decltype(properties(device_image_scope))> Foo;
 
-    void bar(queue q) {
-      q.single_task([=] {
-      Foo = 42;
-    });
-    }
+      void bar(queue q) {
+        q.single_task([=] {
+        Foo = 42;
+      });
+      }
 
-    int main(){ return 0; }
-    "
-    KOKKOS_IMPL_SYCL_DEVICE_GLOBAL_SUPPORTED)
+      int main(){ return 0; }
+      "
+      KOKKOS_IMPL_SYCL_DEVICE_GLOBAL_SUPPORTED)
 
-  IF(KOKKOS_IMPL_SYCL_DEVICE_GLOBAL_SUPPORTED)
-    COMPILER_SPECIFIC_FLAGS(
-      DEFAULT -fsycl-device-code-split=off -DDESUL_SYCL_DEVICE_GLOBAL_SUPPORTED
-    )
+    IF(KOKKOS_IMPL_SYCL_DEVICE_GLOBAL_SUPPORTED)
+      COMPILER_SPECIFIC_FLAGS(
+        DEFAULT -fsycl-device-code-split=off -DDESUL_SYCL_DEVICE_GLOBAL_SUPPORTED
+      )
+    ENDIF()
   ENDIF()
 ENDIF()
 

diff --git a/tpls/desul/include/desul/atomics/Adapt_SYCL.hpp b/tpls/desul/include/desul/atomics/Adapt_SYCL.hpp
@@ -88,15 +88,18 @@ using sycl_atomic_ref = sycl::atomic_ref<T,
                                          sycl::access::address_space::generic_space>;
 #endif
 
-// FIXME_SYCL Use SYCL_EXT_ONEAPI_DEVICE_GLOBAL when available instead
 #ifdef DESUL_SYCL_DEVICE_GLOBAL_SUPPORTED
-// FIXME_SYCL The compiler forces us to use device_image_scope. Drop this when possible.
+#ifdef SYCL_EXT_ONEAPI_DEVICE_GLOBAL
+template <class T>
+using sycl_device_global = sycl::ext::oneapi::experimental::device_global<T>;
+#else
 template <class T>
 using sycl_device_global = sycl::ext::oneapi::experimental::device_global<
     T,
     decltype(sycl::ext::oneapi::experimental::properties(
         sycl::ext::oneapi::experimental::device_image_scope))>;
 #endif
+#endif
 
 }  // namespace Impl
 }  // namespace desul