SYCL Compiler Support, main branch (2024.11.05.) (#301)

* Switched to using <sycl/sycl.hpp>.

As oneAPI 2025.0.0 treats that as the canonical include, giving
warnings about using <CL/sycl.hpp>.

* Introduced explicit support for AdaptiveCpp 24.06.0.

* Removed the last mentions of <CL/sycl.hpp> and cl::sycl.

CMakeLists.txt

@@ -32,6 +32,8 @@ set( CMAKE_ARCHIVE_OUTPUT_DIRECTORY
 # Flags controlling the meta-build system.
 option( VECMEM_USE_SYSTEM_LIBS "Use system libraries by default" FALSE )
 option( VECMEM_BUILD_TESTING "Build the unit tests of VecMem" TRUE )
+option( VECMEM_TEST_UBSAN "Use the undefined behavior sanitizer for the tests"
+   TRUE )
 
 # Include the VecMem CMake code.
 list( APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake" )

cmake/sycl/CMakeDetermineSYCLCompiler.cmake

@@ -24,16 +24,15 @@ if( NOT "$ENV{SYCLCXX}" STREQUAL "" )
    endif()
 
    # Determine the type and version of the SYCL compiler.
-   execute_process( COMMAND "${CMAKE_SYCL_COMPILER_INIT}" "--version"
-      OUTPUT_VARIABLE _syclVersionOutput
-      ERROR_VARIABLE _syclVersionError
-      RESULT_VARIABLE _syclVersionResult )
-   if( NOT ${_syclVersionResult} EQUAL 0 )
-      execute_process( COMMAND "${CMAKE_SYCL_COMPILER_INIT}"
-                               "--hipsycl-version"
+   foreach( _version_cmdl "--acpp-version" "--version" "--hipsycl-version" )
+      execute_process( COMMAND "${CMAKE_SYCL_COMPILER_INIT}" "${_version_cmdl}"
          OUTPUT_VARIABLE _syclVersionOutput
+         ERROR_VARIABLE _syclVersionError
          RESULT_VARIABLE _syclVersionResult )
-   endif()
+      if( ${_syclVersionResult} EQUAL 0 )
+         break()
+      endif()
+   endforeach()
    if( ${_syclVersionResult} EQUAL 0 )
       if( "${_syclVersionOutput}" MATCHES "ComputeCpp" )
          set( CMAKE_SYCL_COMPILER_ID "ComputeCpp" CACHE STRING
@@ -47,6 +46,10 @@ if( NOT "$ENV{SYCLCXX}" STREQUAL "" )
          set( CMAKE_SYCL_COMPILER_ID "IntelLLVM" CACHE STRING
             "Identifier for the SYCL compiler in use" )
          set( _syclVersionRegex "clang version ([0-9\.]+)" )
+      elseif( "${_syclVersionOutput}" MATCHES "AdaptiveCpp" )
+         set( CMAKE_SYCL_COMPILER_ID "AdaptiveCpp" CACHE STRING
+            "Identifier for the SYCL compiler in use" )
+         set( _syclVersionRegex "AdaptiveCpp version: ([0-9\.]+)" )
       elseif( "${_syclVersionOutput}" MATCHES "hipSYCL" )
          set( CMAKE_SYCL_COMPILER_ID "hipSYCL" CACHE STRING
             "Identifier for the SYCL compiler in use" )

cmake/sycl/CMakeTestSYCLCompiler.cmake

@@ -1,6 +1,6 @@
 # VecMem project, part of the ACTS project (R&D line)
 #
-# (c) 2021 CERN for the benefit of the ACTS project
+# (c) 2021-2024 CERN for the benefit of the ACTS project
 #
 # Mozilla Public License Version 2.0
 
@@ -17,7 +17,7 @@ endif()
 # Try to use the HIP compiler.
 file( WRITE
    "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/main.sycl"
-   "#include <CL/sycl.hpp>\n"
+   "#include <sycl/sycl.hpp>\n"
    "int main() {\n"
    "#if (!defined(CL_SYCL_LANGUAGE_VERSION)) &&"
    "    (!defined(SYCL_LANGUAGE_VERSION))\n"

cmake/sycl/Platform/Linux-AdaptiveCpp-SYCL.cmake

@@ -0,0 +1,31 @@
+# VecMem project, part of the ACTS project (R&D line)
+#
+# (c) 2022-2024 CERN for the benefit of the ACTS project
+#
+# Mozilla Public License Version 2.0
+
+# Use the standard GNU compiler options for hipSYCL.
+include( Platform/Linux-GNU )
+__linux_compiler_gnu( SYCL )
+include( Compiler/GNU )
+__compiler_gnu( SYCL )
+
+# Set up the dependency file generation for this platform. Note that SYCL
+# compilation only works with Makefile and Ninja generators, so no check is made
+# here for the current generator.
+set( CMAKE_SYCL_DEPENDS_USE_COMPILER TRUE )
+set( CMAKE_SYCL_DEPFILE_FORMAT gcc )
+
+# Set an archive (static library) creation command explicitly for this platform.
+set( CMAKE_SYCL_CREATE_STATIC_LIBRARY
+   "<CMAKE_AR> qc <TARGET> <LINK_FLAGS> <OBJECTS>" )
+
+# Set the flags controlling the C++ standard used by the SYCL compiler.
+set( CMAKE_SYCL17_STANDARD_COMPILE_OPTION "-std=c++17" )
+set( CMAKE_SYCL17_EXTENSION_COMPILE_OPTION "-std=c++17" )
+
+set( CMAKE_SYCL20_STANDARD_COMPILE_OPTION "-std=c++20" )
+set( CMAKE_SYCL20_EXTENSION_COMPILE_OPTION "-std=c++20" )
+
+set( CMAKE_SYCL23_STANDARD_COMPILE_OPTION "-std=c++23" )
+set( CMAKE_SYCL23_EXTENSION_COMPILE_OPTION "-std=c++23" )

core/cmake/vecmem-setup-core.cmake

@@ -44,39 +44,39 @@ function( vecmem_setup_core libName )
 
       # Test which SYCL printf function(s) is/are available.
       vecmem_check_sycl_source_compiles( "
-         #include <CL/sycl.hpp>
+         #include <sycl/sycl.hpp>
          #ifdef __SYCL_DEVICE_ONLY__
          #  define VECMEM_MSG_ATTRIBUTES __attribute__((opencl_constant))
          #else
          #  define VECMEM_MSG_ATTRIBUTES
          #endif
          int main() {
              const VECMEM_MSG_ATTRIBUTES char __msg[] = \"Test message %i\";
-             cl::sycl::ext::oneapi::experimental::printf(__msg, 20);
+             ::sycl::ext::oneapi::experimental::printf(__msg, 20);
              return 0;
          }
          " VECMEM_HAVE_SYCL_EXT_ONEAPI_PRINTF )
       vecmem_check_sycl_source_compiles( "
-         #include <CL/sycl.hpp>
+         #include <sycl/sycl.hpp>
          #ifdef __SYCL_DEVICE_ONLY__
          #  define VECMEM_MSG_ATTRIBUTES __attribute__((opencl_constant))
          #else
          #  define VECMEM_MSG_ATTRIBUTES
          #endif
          int main() {
              const VECMEM_MSG_ATTRIBUTES char __msg[] = \"Test message %i\";
-             cl::sycl::ONEAPI::experimental::printf(__msg, 20);
+             ::sycl::ONEAPI::experimental::printf(__msg, 20);
              return 0;
          }
          " VECMEM_HAVE_SYCL_ONEAPI_PRINTF )
 
       # Set up the appropriate flag based on these checks.
       if( VECMEM_HAVE_SYCL_EXT_ONEAPI_PRINTF )
          target_compile_definitions( ${libName} INTERFACE
-            $<BUILD_INTERFACE:VECMEM_SYCL_PRINTF_FUNCTION=cl::sycl::ext::oneapi::experimental::printf> )
+            $<BUILD_INTERFACE:VECMEM_SYCL_PRINTF_FUNCTION=::sycl::ext::oneapi::experimental::printf> )
       elseif( VECMEM_HAVE_SYCL_ONEAPI_PRINTF )
          target_compile_definitions( ${libName} INTERFACE
-            $<BUILD_INTERFACE:VECMEM_SYCL_PRINTF_FUNCTION=cl::sycl::ONEAPI::experimental::printf> )
+            $<BUILD_INTERFACE:VECMEM_SYCL_PRINTF_FUNCTION=::sycl::ONEAPI::experimental::printf> )
       else()
          message( WARNING "No valid printf function found for SYCL."
             " Enabling debug messages will likely not work in device code." )
@@ -87,12 +87,12 @@ function( vecmem_setup_core libName )
 
       # Test whether sycl::atomic_ref is available.
       vecmem_check_sycl_source_compiles( "
-         #include <CL/sycl.hpp>
+         #include <sycl/sycl.hpp>
          int main() {
              int dummy = 0;
-             cl::sycl::atomic_ref<int, sycl::memory_order::relaxed,
-                                 cl::sycl::memory_scope::device,
-                                 cl::sycl::access::address_space::global_space>
+             ::sycl::atomic_ref<int, sycl::memory_order::relaxed,
+                               ::sycl::memory_scope::device,
+                               ::sycl::access::address_space::global_space>
                  atomic_dummy(dummy);
              atomic_dummy.store(3);
              atomic_dummy.fetch_add(1);

core/include/vecmem/containers/impl/device_vector.ipp

@@ -1,6 +1,6 @@
 /* VecMem project, part of the ACTS project (R&D line)
  *
- * (c) 2021-2022 CERN for the benefit of the ACTS project
+ * (c) 2021-2024 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
@@ -197,7 +197,7 @@ VECMEM_HOST_AND_DEVICE auto device_vector<TYPE>::emplace_back(Args&&... args)
     // Increment the size of the vector at first. So that we would "claim" the
     // index from other threads.
     device_atomic_ref<size_type> asize(*m_size);
-    const size_type index = asize.fetch_add(1);
+    const size_type index = asize.fetch_add(1u);
     assert(index < m_capacity);
 
     // Instantiate the new value.
@@ -217,7 +217,7 @@ VECMEM_HOST_AND_DEVICE auto device_vector<TYPE>::push_back(
     // Increment the size of the vector at first. So that we would "claim" the
     // index from other threads.
     device_atomic_ref<size_type> asize(*m_size);
-    const size_type index = asize.fetch_add(1);
+    const size_type index = asize.fetch_add(1u);
     assert(index < m_capacity);
 
     // Instantiate the new value.
@@ -315,7 +315,7 @@ VECMEM_HOST_AND_DEVICE auto device_vector<TYPE>::pop_back() -> size_type {
 
     // Decrement the size of the vector, and remember this new size.
     device_atomic_ref<size_type> asize(*m_size);
-    const size_type new_size = asize.fetch_sub(1) - 1;
+    const size_type new_size = asize.fetch_sub(1u) - 1;
 
     // Remove the last element.
     destruct(new_size);

core/include/vecmem/edm/impl/device.ipp

@@ -73,7 +73,7 @@ device<schema<VARTYPES...>, INTERFACE>::push_back_default() -> size_type {
     // Increment the size of the container at first. So that we would "claim"
     // the index from other threads.
     device_atomic_ref<size_type> asize(*m_size);
-    const size_type index = asize.fetch_add(1);
+    const size_type index = asize.fetch_add(1u);
     assert(index < m_capacity);
 
     // Construct the new elements in all of the vector variables.

core/include/vecmem/memory/details/sycl_builtin_device_atomic_ref.hpp

@@ -11,7 +11,7 @@
 #include "vecmem/memory/device_address_space.hpp"
 
 // SYCL include(s).
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 namespace vecmem {
 namespace sycl {
@@ -27,23 +27,20 @@ struct builtin_address_space {};
 /// Specialization for global device memory
 template <>
 struct builtin_address_space<device_address_space::global> {
-    static constexpr cl::sycl::memory_order ord =
-        cl::sycl::memory_order::relaxed;
-    static constexpr cl::sycl::memory_scope scp =
-        cl::sycl::memory_scope::device;
-    static constexpr cl::sycl::access::address_space add =
-        cl::sycl::access::address_space::global_space;
+    static constexpr ::sycl::memory_order ord = ::sycl::memory_order::relaxed;
+    static constexpr ::sycl::memory_scope scp = ::sycl::memory_scope::device;
+    static constexpr ::sycl::access::address_space add =
+        ::sycl::access::address_space::global_space;
 };
 
 /// Specialization for local device memory
 template <>
 struct builtin_address_space<device_address_space::local> {
-    static constexpr cl::sycl::memory_order ord =
-        cl::sycl::memory_order::relaxed;
-    static constexpr cl::sycl::memory_scope scp =
-        cl::sycl::memory_scope::work_group;
-    static constexpr cl::sycl::access::address_space add =
-        cl::sycl::access::address_space::local_space;
+    static constexpr ::sycl::memory_order ord = ::sycl::memory_order::relaxed;
+    static constexpr ::sycl::memory_scope scp =
+        ::sycl::memory_scope::work_group;
+    static constexpr ::sycl::access::address_space add =
+        ::sycl::access::address_space::local_space;
 };
 
 }  // namespace details
@@ -52,9 +49,9 @@ struct builtin_address_space<device_address_space::local> {
 template <typename T,
           device_address_space address = device_address_space::global>
 using builtin_device_atomic_ref =
-    cl::sycl::atomic_ref<T, details::builtin_address_space<address>::ord,
-                         details::builtin_address_space<address>::scp,
-                         details::builtin_address_space<address>::add>;
+    ::sycl::atomic_ref<T, details::builtin_address_space<address>::ord,
+                       details::builtin_address_space<address>::scp,
+                       details::builtin_address_space<address>::add>;
 
 }  // namespace sycl
 }  // namespace vecmem

core/include/vecmem/memory/impl/atomic.ipp

@@ -14,22 +14,21 @@
 
 // SYCL include(s).
 #if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 #endif
 
 /// Helpers for explicit calls to the SYCL atomic functions
 #if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
 #define __VECMEM_SYCL_ATOMIC_CALL0(FNAME, PTR) \
-    cl::sycl::atomic_##FNAME<value_type>(      \
-        cl::sycl::atomic<value_type>(cl::sycl::global_ptr<value_type>(PTR)))
-#define __VECMEM_SYCL_ATOMIC_CALL1(FNAME, PTR, ARG1)                         \
-    cl::sycl::atomic_##FNAME<value_type>(                                    \
-        cl::sycl::atomic<value_type>(cl::sycl::global_ptr<value_type>(PTR)), \
-        ARG1)
-#define __VECMEM_SYCL_ATOMIC_CALL2(FNAME, PTR, ARG1, ARG2)                   \
-    cl::sycl::atomic_##FNAME<value_type>(                                    \
-        cl::sycl::atomic<value_type>(cl::sycl::global_ptr<value_type>(PTR)), \
-        ARG1, ARG2)
+    ::sycl::atomic_##FNAME<value_type>(        \
+        ::sycl::atomic<value_type>(::sycl::global_ptr<value_type>(PTR)))
+#define __VECMEM_SYCL_ATOMIC_CALL1(FNAME, PTR, ARG1) \
+    ::sycl::atomic_##FNAME<value_type>(              \
+        ::sycl::atomic<value_type>(::sycl::global_ptr<value_type>(PTR)), ARG1)
+#define __VECMEM_SYCL_ATOMIC_CALL2(FNAME, PTR, ARG1, ARG2)                     \
+    ::sycl::atomic_##FNAME<value_type>(                                        \
+        ::sycl::atomic<value_type>(::sycl::global_ptr<value_type>(PTR)), ARG1, \
+        ARG2)
 #endif
 
 namespace vecmem {

core/include/vecmem/memory/impl/sycl_custom_device_atomic_ref.ipp

@@ -8,7 +8,7 @@
 #pragma once
 
 // SYCL include(s).
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 namespace vecmem {
 namespace sycl {
@@ -19,43 +19,43 @@ struct custom_address_space {};
 
 template <>
 struct custom_address_space<device_address_space::global> {
-    static constexpr cl::sycl::access::address_space add =
-        cl::sycl::access::address_space::global_space;
+    static constexpr ::sycl::access::address_space add =
+        ::sycl::access::address_space::global_space;
 
     template <typename T>
-    using ptr_t = cl::sycl::global_ptr<T>;
+    using ptr_t = ::sycl::global_ptr<T>;
 };
 
 template <>
 struct custom_address_space<device_address_space::local> {
-    static constexpr cl::sycl::access::address_space add =
-        cl::sycl::access::address_space::local_space;
+    static constexpr ::sycl::access::address_space add =
+        ::sycl::access::address_space::local_space;
     template <typename T>
-    using ptr_t = cl::sycl::local_ptr<T>;
+    using ptr_t = ::sycl::local_ptr<T>;
 };
 
 }  // namespace details
 
 #define __VECMEM_SYCL_ATOMIC_CALL0(FNAME, PTR)                               \
-    cl::sycl::atomic_##FNAME<value_type,                                     \
-                             details::custom_address_space<address>::add>(   \
-        cl::sycl::atomic<value_type,                                         \
-                         details::custom_address_space<address>::add>(       \
+    ::sycl::atomic_##FNAME<value_type,                                       \
+                           details::custom_address_space<address>::add>(     \
+        ::sycl::atomic<value_type,                                           \
+                       details::custom_address_space<address>::add>(         \
             typename details::custom_address_space<address>::template ptr_t< \
                 value_type>(PTR)))
 #define __VECMEM_SYCL_ATOMIC_CALL1(FNAME, PTR, ARG1)                         \
-    cl::sycl::atomic_##FNAME<value_type,                                     \
-                             details::custom_address_space<address>::add>(   \
-        cl::sycl::atomic<value_type,                                         \
-                         details::custom_address_space<address>::add>(       \
+    ::sycl::atomic_##FNAME<value_type,                                       \
+                           details::custom_address_space<address>::add>(     \
+        ::sycl::atomic<value_type,                                           \
+                       details::custom_address_space<address>::add>(         \
             typename details::custom_address_space<address>::template ptr_t< \
                 value_type>(PTR)),                                           \
         ARG1)
 #define __VECMEM_SYCL_ATOMIC_CALL2(FNAME, PTR, ARG1, ARG2)                   \
-    cl::sycl::atomic_##FNAME<value_type,                                     \
-                             details::custom_address_space<address>::add>(   \
-        cl::sycl::atomic<value_type,                                         \
-                         details::custom_address_space<address>::add>(       \
+    ::sycl::atomic_##FNAME<value_type,                                       \
+                           details::custom_address_space<address>::add>(     \
+        ::sycl::atomic<value_type,                                           \
+                       details::custom_address_space<address>::add>(         \
             typename details::custom_address_space<address>::template ptr_t< \
                 value_type>(PTR)),                                           \
         ARG1, ARG2)

core/include/vecmem/memory/memory_order.hpp

@@ -14,7 +14,7 @@
 #if (defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)) && \
     defined(VECMEM_HAVE_SYCL_ATOMIC_REF)
 // SYCL include(s).
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 #endif
 
 namespace vecmem {

sycl/CMakeLists.txt

@@ -41,7 +41,7 @@ set_target_properties( vecmem_sycl PROPERTIES
    CXX_VISIBILITY_PRESET  "hidden"
    SYCL_VISIBILITY_PRESET "hidden" )
 vecmem_check_sycl_source_compiles( "
-   #include <CL/sycl.hpp>
+   #include <sycl/sycl.hpp>
    int main() { return 0; }
    "
    VECMEM_HAVE_SYCL_VISIBILITY_MS_COMPAT
@@ -94,9 +94,9 @@ endif()
 # Check if sycl::queue::memset is available, and set a compiler option
 # accordingly.
 vecmem_check_sycl_source_compiles( "
-   #include <CL/sycl.hpp>
+   #include <sycl/sycl.hpp>
    int main() {
-       cl::sycl::queue queue;
+       ::sycl::queue queue;
        queue.memset(nullptr, 0, 100);
        return 0;
    }