oneAPI 2025 / SYCL2020 Compatibility, main branch (2024.12.03.)

.github/ci_setup.sh

@@ -19,7 +19,7 @@ export CTEST_PARALLEL_LEVEL=${CMAKE_BUILD_PARALLEL_LEVEL}
 export MAKEFLAGS="-j${CMAKE_BUILD_PARALLEL_LEVEL}"
 
 # Set up the correct environment for the SYCL tests.
-if [ "${PLATFORM_NAME}" = "SYCL" ]; then
+if [[ "${PLATFORM_NAME}" == *"SYCL"* ]]; then
    if [ -f "/opt/intel/oneapi/setvars.sh" ]; then
       OLD_CPATH=${CPATH}
       source /opt/intel/oneapi/setvars.sh --include-intel-llvm

.github/workflows/builds.yml

@@ -25,23 +25,23 @@ jobs:
       matrix:
         platform:
           - name: CPU
-            container: ghcr.io/acts-project/ubuntu2404:56
+            container: ghcr.io/acts-project/ubuntu2404:69
             options: --preset host-fp32 -DALGEBRA_PLUGINS_USE_SYSTEM_VC=FALSE
             run_tests: true
           - name: CPU
-            container: ghcr.io/acts-project/ubuntu2404:56
+            container: ghcr.io/acts-project/ubuntu2404:69
             options: --preset host-fp64 -DALGEBRA_PLUGINS_USE_SYSTEM_VC=FALSE
             run_tests: false
           - name: CUDA
-            container: ghcr.io/acts-project/ubuntu2404_cuda:56
+            container: ghcr.io/acts-project/ubuntu2404_cuda:69
             options: --preset cuda-fp32
             run_tests: false
-          - name: SYCL
-            container: ghcr.io/acts-project/ubuntu2404_oneapi:56
-            options: --preset sycl-fp32 -DCMAKE_SYCL_FLAGS="-fsycl -fsycl-targets=spir64"
+          - name: "SYCL Intel"
+            container: ghcr.io/acts-project/ubuntu2404_oneapi:69
+            options: --preset sycl-fp32
             run_tests: true
           - name: KOKKOS
-            container: ghcr.io/acts-project/ubuntu2404:56
+            container: ghcr.io/acts-project/ubuntu2404:69
             options: --preset kokkos-fp32
             run_tests: false
         build:
@@ -50,31 +50,31 @@ jobs:
         include:
           - platform:
               name: CUDA
-              container: ghcr.io/acts-project/ubuntu2404_cuda:56
+              container: ghcr.io/acts-project/ubuntu2404_cuda:69
               options: --preset cuda-fp64
               run_tests: false
             build: Release
           - platform:
-              name: "SYCL"
-              container: ghcr.io/acts-project/ubuntu2404_cuda_oneapi:56
+              name: "SYCL NVIDIA"
+              container: ghcr.io/acts-project/ubuntu2404_cuda_oneapi:69
               options: --preset sycl-fp32
               run_tests: false
             build: Release
           - platform:
-              name: "SYCL"
-              container: ghcr.io/acts-project/ubuntu2404_rocm_oneapi:56
+              name: "SYCL AMD"
+              container: ghcr.io/acts-project/ubuntu2404_rocm_oneapi:69
               options: --preset sycl-fp32
               run_tests: false
             build: Release
           - platform:
               name: ALPAKA
-              container: ghcr.io/acts-project/ubuntu2404:56
+              container: ghcr.io/acts-project/ubuntu2404:69
               options: --preset alpaka-fp32
               run_tests: true
             build: Release
           - platform:
               name: ALPAKA
-              container: ghcr.io/acts-project/ubuntu2404:56
+              container: ghcr.io/acts-project/ubuntu2404:69
               options: --preset alpaka-fp32
               run_tests: false
             build: Debug
@@ -83,7 +83,13 @@ jobs:
       run:
         shell: bash
     steps:
-      - uses: actions/checkout@v3
+      - uses: actions/checkout@v4
+      - name: Install dependencies
+        run: |
+          apt install -y zstd
+          curl --retry 5 --retry-delay 10 --output deps.tar.zst https://acts.web.cern.ch/ACTS/ci/ubuntu-24.04/deps.v6.tar.zst
+          tar -xf deps.tar.zst -C /usr/local --strip-components=1
+          rm deps.tar.zst
       - name: Configure
         run: |
           source ${GITHUB_WORKSPACE}/.github/ci_setup.sh ${{ matrix.platform.name }}

core/include/traccc/definitions/math.hpp

@@ -9,7 +9,7 @@
 
 // SYCL include(s).
 #if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 #endif
 
 // System include(s).
@@ -19,7 +19,7 @@ namespace traccc {
 
 /// Namespace to pick up math functions from
 #if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
-namespace math = cl::sycl;
+namespace math = ::sycl;
 #else
 namespace math = std;
 #endif  // SYCL

device/sycl/include/traccc/sycl/utils/calculate1DimNdRange.hpp

@@ -1,14 +1,14 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2022 CERN for the benefit of the ACTS project
+ * (c) 2022-2024 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
 
 #pragma once
 
 // Sycl include(s).
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 namespace traccc::sycl {
 

device/sycl/include/traccc/sycl/utils/make_prefix_sum_buff.hpp

@@ -1,14 +1,14 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2022 CERN for the benefit of the ACTS project
+ * (c) 2022-2024 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
 
 #pragma once
 
 // Sycl include(s).
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 // Project include(s).
 #include "traccc/device/fill_prefix_sum.hpp"

device/sycl/src/fitting/fit_tracks.hpp

@@ -27,7 +27,7 @@
 #include <oneapi/dpl/execution>
 
 // SYCL include(s).
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 namespace traccc::sycl {
 namespace kernels {
@@ -46,7 +46,7 @@ track_state_container_types::buffer fit_tracks(
     const typename track_candidate_container_types::const_view&
         track_candidates_view,
     const fitting_config& config, const memory_resource& mr, vecmem::copy& copy,
-    cl::sycl::queue& queue) {
+    ::sycl::queue& queue) {
 
     // Get the number of tracks.
     const track_candidate_container_types::const_device::header_vector::
@@ -87,15 +87,15 @@ track_state_container_types::buffer fit_tracks(
 
     // The execution range for the two kernels of the function.
     static constexpr unsigned int localSize = 64;
-    cl::sycl::nd_range<1> range = calculate1DimNdRange(n_tracks, localSize);
+    ::sycl::nd_range<1> range = calculate1DimNdRange(n_tracks, localSize);
 
     // Fill the keys and param_ids buffers.
-    cl::sycl::event fill_keys_event = queue.submit([&](cl::sycl::handler& h) {
+    ::sycl::event fill_keys_event = queue.submit([&](::sycl::handler& h) {
         h.parallel_for<kernels::fill_sort_keys>(
             range,
             [track_candidates_view, keys_view = vecmem::get_data(keys_buffer),
-             param_ids_view = vecmem::get_data(param_ids_buffer)](
-                cl::sycl::nd_item<1> item) {
+             param_ids_view =
+                 vecmem::get_data(param_ids_buffer)](::sycl::nd_item<1> item) {
                 device::fill_sort_keys(item.get_global_linear_id(),
                                        track_candidates_view, keys_view,
                                        param_ids_view);
@@ -115,11 +115,11 @@ track_state_container_types::buffer fit_tracks(
     track_states_headers_setup_event->wait();
     track_states_items_setup_event->wait();
     queue
-        .submit([&](cl::sycl::handler& h) {
+        .submit([&](::sycl::handler& h) {
             h.parallel_for<fit_kernel_t>(
                 range, [det_view, field_view, config, track_candidates_view,
                         param_ids_view = vecmem::get_data(param_ids_buffer),
-                        track_states_view](cl::sycl::nd_item<1> item) {
+                        track_states_view](::sycl::nd_item<1> item) {
                     device::fit<fitter_t>(item.get_global_linear_id(), det_view,
                                           field_view, config,
                                           track_candidates_view, param_ids_view,

device/sycl/src/sanity/contiguous_on.hpp

@@ -19,7 +19,7 @@
 #include <vecmem/utils/copy.hpp>
 
 // SYCL include
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 // System include
 #include <concepts>
@@ -44,7 +44,7 @@ struct is_contiguous_on_compress_adjacent {
         : m_projection(projection), m_view(view), m_out_view(out_view) {}
 
     /// Execution operator for the kernel
-    void operator()(cl::sycl::nd_item<1> item) const {
+    void operator()(::sycl::nd_item<1> item) const {
 
         std::size_t tid = item.get_global_linear_id();
 
@@ -111,7 +111,7 @@ is_contiguous_on(P&& projection, vecmem::memory_resource& mr,
     constexpr int local_size = 512;
     constexpr int local_size_2d = 32;
 
-    cl::sycl::queue& queue = details::get_queue(queue_wrapper);
+    ::sycl::queue& queue = details::get_queue(queue_wrapper);
 
     // Grab the number of elements in our vector.
     const typename VIEW::size_type n = copy.get_size(view);
@@ -133,16 +133,16 @@ is_contiguous_on(P&& projection, vecmem::memory_resource& mr,
 
     bool initial_out = true;
 
-    cl::sycl::event kernel2_memcpy_evt = queue.copy(&initial_out, out.get(), 1);
+    ::sycl::event kernel2_memcpy_evt = queue.copy(&initial_out, out.get(), 1);
 
-    cl::sycl::nd_range<1> compress_adjacent_range{
-        cl::sycl::range<1>(((n + local_size - 1) / local_size) * local_size),
-        cl::sycl::range<1>(local_size)};
+    ::sycl::nd_range<1> compress_adjacent_range{
+        ::sycl::range<1>(((n + local_size - 1) / local_size) * local_size),
+        ::sycl::range<1>(local_size)};
 
     // Launch the first kernel, which will squash consecutive equal elements
     // into one element.
     queue
-        .submit([&](cl::sycl::handler& h) {
+        .submit([&](::sycl::handler& h) {
             h.parallel_for<kernels::is_contiguous_on_compress_adjacent<
                 CONTAINER, P, VIEW, projection_t>>(
                 compress_adjacent_range,
@@ -156,17 +156,17 @@ is_contiguous_on(P&& projection, vecmem::memory_resource& mr,
         copy.get_size(iout);
     uint32_t grid_size_rd =
         (host_iout_size + local_size_2d - 1) / local_size_2d;
-    cl::sycl::nd_range<2> all_unique_range{
-        cl::sycl::range<2>(grid_size_rd * local_size_2d,
-                           grid_size_rd * local_size_2d),
-        cl::sycl::range<2>(local_size_2d, local_size_2d)};
+    ::sycl::nd_range<2> all_unique_range{
+        ::sycl::range<2>(grid_size_rd * local_size_2d,
+                         grid_size_rd * local_size_2d),
+        ::sycl::range<2>(local_size_2d, local_size_2d)};
 
     // Launch the second kernel, which will check if the values are unique.
-    cl::sycl::event kernel2_evt = queue.submit([&](cl::sycl::handler& h) {
+    ::sycl::event kernel2_evt = queue.submit([&](::sycl::handler& h) {
         h.depends_on(kernel2_memcpy_evt);
         h.parallel_for<kernels::is_contiguous_on_all_unique<projection_t>>(
             all_unique_range, [in_view = vecmem::get_data(iout),
-                               out = out.get()](cl::sycl::nd_item<2> item) {
+                               out = out.get()](::sycl::nd_item<2> item) {
                 std::size_t tid_x = item.get_global_id(0);
                 std::size_t tid_y = item.get_global_id(1);
 

device/sycl/src/sanity/ordered_on.hpp

@@ -18,7 +18,7 @@
 #include <vecmem/utils/copy.hpp>
 
 // SYCL include
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 // System include
 #include <concepts>
@@ -40,7 +40,7 @@ struct is_ordered_on {
         : m_relation(relation), m_view(view), m_out(out) {}
 
     /// Execution operator for the kernel
-    void operator()(cl::sycl::nd_item<1> item) const {
+    void operator()(::sycl::nd_item<1> item) const {
 
         std::size_t tid = item.get_global_linear_id();
 
@@ -98,7 +98,7 @@ is_ordered_on(R&& relation, vecmem::memory_resource& mr, vecmem::copy& copy,
     // block size fixed.
     constexpr int block_size = 512;
 
-    cl::sycl::queue& queue = details::get_queue(queue_wrapper);
+    ::sycl::queue& queue = details::get_queue(queue_wrapper);
 
     // Grab the number of elements in our container.
     const typename VIEW::size_type n = copy.get_size(view);
@@ -112,14 +112,14 @@ is_ordered_on(R&& relation, vecmem::memory_resource& mr, vecmem::copy& copy,
     vecmem::unique_alloc_ptr<bool> out = vecmem::make_unique_alloc<bool>(mr);
     bool initial_out = true;
 
-    cl::sycl::event kernel1_memcpy1 =
+    ::sycl::event kernel1_memcpy1 =
         queue.memcpy(out.get(), &initial_out, sizeof(bool));
 
-    cl::sycl::nd_range<1> kernel_range{
-        cl::sycl::range<1>(((n + block_size - 1) / block_size) * block_size),
-        cl::sycl::range<1>(block_size)};
+    ::sycl::nd_range<1> kernel_range{
+        ::sycl::range<1>(((n + block_size - 1) / block_size) * block_size),
+        ::sycl::range<1>(block_size)};
 
-    cl::sycl::event kernel1 = queue.submit([&](cl::sycl::handler& h) {
+    ::sycl::event kernel1 = queue.submit([&](::sycl::handler& h) {
         h.depends_on(kernel1_memcpy1);
         h.parallel_for<kernels::is_ordered_on<CONTAINER, R, VIEW>>(
             kernel_range, kernels::is_ordered_on<CONTAINER, R, VIEW>(

device/sycl/src/seeding/silicon_pixel_spacepoint_formation.hpp

@@ -20,7 +20,7 @@
 #include <vecmem/memory/memory_resource.hpp>
 
 // SYCL include(s).
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 namespace traccc::sycl::details {
 
@@ -40,7 +40,7 @@ template <typename detector_t>
 spacepoint_collection_types::buffer silicon_pixel_spacepoint_formation(
     const typename detector_t::view_type& det_view,
     const measurement_collection_types::const_view& measurements_view,
-    vecmem::memory_resource& mr, vecmem::copy& copy, cl::sycl::queue& queue) {
+    vecmem::memory_resource& mr, vecmem::copy& copy, ::sycl::queue& queue) {
 
     // Get the number of measurements.
     const measurement_collection_types::const_view::size_type n_measurements =
@@ -63,11 +63,11 @@ spacepoint_collection_types::buffer silicon_pixel_spacepoint_formation(
 
     // Run the spacepoint formation on the device.
     queue
-        .submit([&](cl::sycl::handler& h) {
+        .submit([&](::sycl::handler& h) {
             h.parallel_for(
                 countRange, [det_view, measurements_view, n_measurements,
                              spacepoints_view = vecmem::get_data(result)](
-                                cl::sycl::nd_item<1> item) {
+                                ::sycl::nd_item<1> item) {
                     device::form_spacepoints<detector_t>(
                         item.get_global_linear_id(), det_view,
                         measurements_view, n_measurements, spacepoints_view);

device/sycl/src/seeding/spacepoint_binning.sycl

@@ -17,7 +17,7 @@
 #include "traccc/seeding/device/populate_grid.hpp"
 
 // SYCL include(s).
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 namespace traccc::sycl {
 namespace kernels {

device/sycl/src/utils/barrier.hpp

@@ -11,12 +11,12 @@
 #include "traccc/definitions/qualifiers.hpp"
 
 // SYCL includes
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 namespace traccc::sycl {
 
 struct barrier {
-    barrier(::sycl::nd_item<1> item) : m_item(item){};
+    barrier(::sycl::nd_item<1> item) : m_item(item) {}
 
     TRACCC_DEVICE
     void blockBarrier() { m_item.barrier(); }

device/sycl/src/utils/get_queue.hpp

@@ -11,7 +11,7 @@
 #include "traccc/sycl/utils/queue_wrapper.hpp"
 
 // SYCL include(s).
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 namespace traccc::sycl::details {
 

examples/run/sycl/CMakeLists.txt

@@ -13,10 +13,6 @@ include( traccc-compiler-options-sycl )
 # We need the SYCL language in this directory.
 enable_language( SYCL )
 
-# Example executable(s).
-traccc_add_executable( traccc_sycl_language_example
-   "sycl_language_example.sycl" )
-
 # SYCL seeding executable(s).
 traccc_add_executable( seeding_example_sycl "seeding_example_sycl.sycl"
    LINK_LIBRARIES traccc::options vecmem::core vecmem::sycl traccc::io