Alpaka inhomogeneous magnetic field

.github/workflows/builds.yml

@@ -71,7 +71,7 @@ jobs:
           - platform:
               name: "ALPAKA_CUDA"
               container: ghcr.io/acts-project/ubuntu2404_cuda:69
-              options: --preset alpaka-fp32 -Dalpaka_ACC_GPU_CUDA_ENABLE=ON -DTRACCC_FAIL_ON_WARNINGS=OFF
+              options: --preset alpaka-fp32 -Dalpaka_ACC_GPU_CUDA_ENABLE=ON -DTRACCC_FAIL_ON_WARNINGS=OFF -DTRACCC_BUILD_CUDA=ON
               run_tests: false
             build: Release
           - platform:
@@ -83,7 +83,7 @@ jobs:
           - platform:
               name: "ALPAKA_SYCL"
               container: ghcr.io/acts-project/ubuntu2404_oneapi:69
-              options: --preset alpaka-fp32 -Dalpaka_ACC_CPU_B_SEQ_T_THREADS_ENABLE=OFF -Dalpaka_ACC_SYCL_ENABLE=ON -Dalpaka_SYCL_ONEAPI_GPU=ON -Dalpaka_SYCL_ONEAPI_GPU_DEVICES=spir64
+              options: --preset alpaka-fp32 -Dalpaka_ACC_CPU_B_SEQ_T_THREADS_ENABLE=OFF -Dalpaka_ACC_SYCL_ENABLE=ON -Dalpaka_SYCL_ONEAPI_GPU=ON -Dalpaka_SYCL_ONEAPI_GPU_DEVICES=spir64 -DTRACCC_BUILD_SYCL=ON
               run_tests: false
             build: Release
     # Use BASH as the shell from the images.

CMakeLists.txt

@@ -58,6 +58,7 @@ set( TRACCC_DEVICE_LOG_LVL "NONE" CACHE STRING
 
 # Flags controlling which parts of traccc to build.
 option( TRACCC_BUILD_CUDA "Build the CUDA sources included in traccc" FALSE )
+option( TRACCC_BUILD_CUDA_UTILS "Build the CUDA util sources included in traccc" FALSE )
 option( TRACCC_BUILD_HIP "Build the HIP sources included in traccc" FALSE)
 option( TRACCC_BUILD_SYCL "Build the SYCL sources included in traccc" FALSE )
 option( TRACCC_BUILD_ALPAKA "Build the Alpaka sources included in traccc"
@@ -425,7 +426,7 @@ set(TRACCC_ALGEBRA_PLUGINS ARRAY CACHE STRING "Algebra plugin to use in the buil
 # Build the traccc code.
 add_subdirectory( core )
 add_subdirectory( device/common )
-if( TRACCC_BUILD_CUDA )
+if( TRACCC_BUILD_CUDA OR  TRACCC_BUILD_CUDA_UTILS)
    add_subdirectory( device/cuda )
 endif()
 if( TRACCC_BUILD_SYCL )

device/alpaka/CMakeLists.txt

@@ -14,20 +14,22 @@ set(PRIVATE_LIBRARIES alpaka::alpaka)
 traccc_enable_language_alpaka()
 
 if(alpaka_ACC_GPU_CUDA_ENABLE)
-  list(APPEND PRIVATE_LIBRARIES CUDA::cudart vecmem::cuda)
+  list(APPEND PRIVATE_LIBRARIES CUDA::cudart vecmem::cuda traccc_cuda_utils covfie::cuda)
 elseif(alpaka_ACC_GPU_HIP_ENABLE)
   find_package( HIPToolkit REQUIRED )
-  list(APPEND PRIVATE_LIBRARIES HIP::hiprt vecmem::hip)
+  list(APPEND PRIVATE_LIBRARIES HIP::hiprt vecmem::hip covfie::hip)
 elseif(alpaka_ACC_SYCL_ENABLE)
   include( traccc-compiler-options-sycl )
-  list(APPEND PRIVATE_LIBRARIES vecmem::sycl)
+  list(APPEND PRIVATE_LIBRARIES vecmem::sycl traccc::sycl covfie::sycl)
 endif()
 
 traccc_add_alpaka_library( traccc_alpaka alpaka TYPE SHARED
   # Utility definitions.
   "include/traccc/alpaka/utils/get_device_info.hpp"
   "include/traccc/alpaka/utils/make_prefix_sum_buff.hpp"
+  "include/traccc/alpaka/utils/make_magnetic_field.hpp"
   "src/utils/make_prefix_sum_buff.cpp"
+  "src/utils/make_magnetic_field.cpp"
   "src/utils/get_device_info.cpp"
   "include/traccc/alpaka/utils/queue.hpp"
   "src/utils/queue.cpp"

device/alpaka/include/traccc/alpaka/utils/make_magnetic_field.hpp

@@ -0,0 +1,23 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Project include(s).
+#include "traccc/alpaka/utils/queue.hpp"
+#include "traccc/bfield/magnetic_field.hpp"
+
+namespace traccc::alpaka {
+
+/// Create a magnetic field usable on the active device
+///
+/// @param bfield The magnetic field to be copied
+//
+magnetic_field make_magnetic_field(const magnetic_field& bfield,
+                                   const queue& queue);
+
+}  // namespace traccc::alpaka

device/alpaka/src/utils/magnetic_field_types.hpp

@@ -8,12 +8,53 @@
 #pragma once
 
 #include "traccc/bfield/magnetic_field_types.hpp"
+#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
+#include <covfie/cuda/backend/primitive/cuda_device_array.hpp>
+#elif defined(ALPAKA_ACC_GPU_HIP_ENABLED)
+#include <covfie/hip/backend/primitive/hip_device_array.hpp>
+#elif defined(ALPAKA_ACC_SYCL_ENABLED)
+#include <covfie/sycl/backend/primitive/sycl_device_array.hpp>
+#endif
 
 namespace traccc::alpaka {
 
+/// Inhomogeneous B-field backend type for Alpaka
+
+#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
+/// Inhomogeneous B-field backend type using CUDA global memory
+template <typename scalar_t>
+using inhom_bfield_backend_t =
+    covfie::backend::affine<covfie::backend::linear<covfie::backend::clamp<
+        covfie::backend::strided<covfie::vector::vector_d<std::size_t, 3>,
+                                 covfie::backend::cuda_device_array<
+                                     covfie::vector::vector_d<scalar_t, 3>>>>>>;
+#elif defined(ALPAKA_ACC_GPU_HIP_ENABLED)
+template <typename scalar_t>
+using inhom_bfield_backend_t =
+    covfie::backend::affine<covfie::backend::linear<covfie::backend::clamp<
+        covfie::backend::strided<covfie::vector::vector_d<std::size_t, 3>,
+                                 covfie::backend::hip_device_array<
+                                     covfie::vector::vector_d<scalar_t, 3>>>>>>;
+#elif defined(ALPAKA_ACC_SYCL_ENABLED)
+template <typename scalar_t>
+using inhom_bfield_backend_t =
+    covfie::backend::affine<covfie::backend::linear<covfie::backend::clamp<
+        covfie::backend::strided<covfie::vector::vector_d<std::size_t, 3>,
+                                 covfie::backend::sycl_device_array<
+                                     covfie::vector::vector_d<scalar_t, 3>>>>>>;
+#else
+template <typename scalar_t>
+using inhom_bfield_backend_t = traccc::host::inhom_bfield_backend_t<scalar_t>;
+#endif
+
+// Test that the type is a valid backend for a field
+static_assert(
+    covfie::concepts::field_backend<inhom_bfield_backend_t<float>>,
+    "alpaka::inhom_bfield_backend_t is not a valid field backend type");
 /// @brief the standard list of Alpaka bfield types to support
 template <typename scalar_t>
 using bfield_type_list = std::tuple<const_bfield_backend_t<scalar_t>,
-                                    host::inhom_bfield_backend_t<scalar_t>>;
+                                    host::inhom_bfield_backend_t<scalar_t>,
+                                    inhom_bfield_backend_t<scalar_t>>;
 
 }  // namespace traccc::alpaka

device/alpaka/src/utils/make_magnetic_field.cpp

@@ -0,0 +1,55 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+// Local include(s).
+#include "traccc/alpaka/utils/make_magnetic_field.hpp"
+
+#include "magnetic_field_types.hpp"
+
+// Project include(s).
+#include "traccc/definitions/primitives.hpp"
+#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
+#include "traccc/cuda/utils/make_magnetic_field.hpp"
+#elif defined(ALPAKA_ACC_GPU_HIP_ENABLED)
+//
+#elif defined(ALPAKA_ACC_SYCL_ENABLED)
+#include "../utils/get_queue.hpp"
+#include "traccc/sycl/utils/make_magnetic_field.hpp"
+#endif
+
+// System include(s).
+#include <stdexcept>
+
+namespace traccc::alpaka {
+
+magnetic_field make_magnetic_field(const magnetic_field& bfield,
+                                   [[maybe_unused]] const queue& queue) {
+#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
+    return traccc::cuda::make_magnetic_field(
+        bfield, traccc::cuda::magnetic_field_storage::global_memory);
+#elif defined(ALPAKA_ACC_GPU_HIP_ENABLED)
+    if (bfield.is<const_bfield_backend_t<traccc::scalar> >()) {
+        return magnetic_field{covfie::field<const_bfield_backend_t<scalar> >{
+            bfield.as_field<const_bfield_backend_t<traccc::scalar> >()}};
+    } else if (bfield.is<host::inhom_bfield_backend_t<traccc::scalar> >()) {
+        return magnetic_field{covfie::field<
+            inhom_bfield_backend_t<traccc::scalar> >(
+            bfield.as_field<host::inhom_bfield_backend_t<traccc::scalar> >())};
+    } else {
+        throw std::invalid_argument(
+            "Unsupported storage method chosen for inhomogeneous b-field");
+    }
+#elif defined(ALPAKA_ACC_SYCL_ENABLED)
+    ::sycl::queue q(::alpaka::getNativeHandle(details::get_queue(queue)));
+    traccc::sycl::queue_wrapper qw{&q};
+    return traccc::sycl::make_magnetic_field(bfield, qw);
+#else
+    return bfield;
+#endif
+}
+
+}  // namespace traccc::alpaka

device/cuda/CMakeLists.txt

@@ -17,8 +17,8 @@ find_package( CUDAToolkit REQUIRED )
 include( traccc-compiler-options-cpp )
 include( traccc-compiler-options-cuda )
 
-# Set up the build of the traccc::cuda library.
-traccc_add_library( traccc_cuda cuda TYPE SHARED
+# Set up the build of the traccc::cuda_utils library.
+traccc_add_library( traccc_cuda_utils cuda TYPE SHARED
   # Utility definitions.
   "include/traccc/cuda/utils/make_prefix_sum_buff.hpp"
   "src/utils/make_prefix_sum_buff.cu"
@@ -36,6 +36,10 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED
   "include/traccc/cuda/utils/make_magnetic_field.hpp"
   "src/utils/make_magnetic_field.cpp"
   "src/utils/magnetic_field_types.hpp"
+)
+# Set up the build of the traccc::cuda library.
+if (TRACCC_BUILD_CUDA)
+traccc_add_library( traccc_cuda cuda TYPE SHARED
   # Seed finding code.
   "include/traccc/cuda/seeding/track_params_estimation.hpp"
   "src/seeding/track_params_estimation.cu"
@@ -117,9 +121,11 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED
   "src/fitting/kernels/fill_fitting_sort_keys.cu"
   "src/fitting/kernels/fit_prelude.cu"
 )
+endif()
 
 set(TRACCC_CUDA_SUPPORTED_BFIELDS "const;inhom_global;inhom_texture")
 
+if (TRACCC_BUILD_CUDA)
 set(KERNEL_SPECIALIZATION_PY "${PROJECT_SOURCE_DIR}/codegen/kernel_specialization/gen_kernel_specialization.py")
 
 # Generate specializations of find_tracks
@@ -228,6 +234,7 @@ foreach(DETECTOR_NAME ${TRACCC_SUPPORTED_DETECTORS})
         target_sources(traccc_cuda PRIVATE ${GENERATED_SOURCE})
     endforeach()
 endforeach()
+endif()
 
 if(TRACCC_ENABLE_NVTX_PROFILING)
     traccc_add_library(
@@ -259,18 +266,26 @@ if(TRACCC_ENABLE_NVTX_PROFILING)
     )
 endif()
 
-target_compile_options( traccc_cuda
+target_compile_options( traccc_cuda_utils
   PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr> )
-target_link_libraries( traccc_cuda
+target_link_libraries( traccc_cuda_utils
   PUBLIC traccc::core detray::core vecmem::core
   PRIVATE CUDA::cudart traccc::device_common vecmem::cuda covfie::cuda )
+if(TRACCC_BUILD_CUDA)
+target_compile_options( traccc_cuda
+  PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr> )
+target_link_libraries( traccc_cuda
+  PRIVATE CUDA::cudart traccc::device_common vecmem::cuda covfie::cuda traccc_cuda_utils )
+endif()
 
 # Set up Thrust specifically for the traccc::cuda library.
 thrust_create_target( traccc::cuda_thrust
   HOST CPP
   DEVICE CUDA )
 target_link_libraries( traccc::cuda_thrust INTERFACE CUDA::cudart )
+if(TRACCC_BUILD_CUDA)
 target_link_libraries( traccc_cuda PRIVATE traccc::cuda_thrust )
+endif()
 
 # For CUDA 11 turn on separable compilation. This is necessary for using
 # Thrust 2.1.0.

device/sycl/src/utils/magnetic_field_types.hpp

@@ -24,7 +24,7 @@
 
 namespace traccc::sycl {
 
-/// Inhomogeneous B-field backend type for CUDA
+/// Inhomogeneous B-field backend type for SYCL
 template <typename scalar_t>
 using inhom_bfield_backend_t =
     covfie::backend::affine<covfie::backend::linear<covfie::backend::clamp<

examples/run/alpaka/full_chain_algorithm.cpp

@@ -8,6 +8,9 @@
 // Local include(s).
 #include "full_chain_algorithm.hpp"
 
+// Project include(s).
+#include "traccc/alpaka/utils/make_magnetic_field.hpp"
+
 // System include(s).
 #include <iostream>
 #include <stdexcept>
@@ -33,7 +36,7 @@ full_chain_algorithm::full_chain_algorithm(
       m_cached_pinned_host_mr(m_vecmem_objects.host_mr()),
       m_cached_device_mr(m_vecmem_objects.device_mr()),
       m_field_vec{0.f, 0.f, finder_config.bFieldInZ},
-      m_field(field),
+      m_field(make_magnetic_field(field, m_queue)),
       m_det_descr(det_descr),
       m_device_det_descr(
           static_cast<silicon_detector_description::buffer::size_type>(

examples/run/cuda/CMakeLists.txt

@@ -12,20 +12,20 @@ find_package( CUDAToolkit REQUIRED )
 
 traccc_add_executable( seq_example_cuda "seq_example_cuda.cpp"
    LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance
-                  traccc::core traccc::device_common traccc::cuda
+                  traccc::core traccc::device_common traccc_cuda_utils traccc::cuda
                   traccc::options detray::detectors detray::io
                   traccc_examples_common )
 traccc_add_executable( seeding_example_cuda "seeding_example_cuda.cpp"
    LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance
-                  traccc::core traccc::device_common traccc::cuda
+                  traccc::core traccc::device_common traccc_cuda_utils traccc::cuda
                   traccc::options traccc_examples_common )
 traccc_add_executable( truth_finding_example_cuda "truth_finding_example_cuda.cpp"
    LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance
-                  traccc::core traccc::device_common traccc::cuda
+                  traccc::core traccc::device_common traccc_cuda_utils traccc::cuda
                   traccc::options traccc_examples_common )
 traccc_add_executable( truth_fitting_example_cuda "truth_fitting_example_cuda.cpp"
    LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance
-                  traccc::core traccc::device_common traccc::cuda
+                  traccc::core traccc::device_common traccc_cuda_utils traccc::cuda
                   traccc::options traccc_examples_common )
 #
 # Set up the "throughput applications".
@@ -35,17 +35,17 @@ add_library( traccc_examples_cuda STATIC
    "full_chain_algorithm.cpp" )
 target_link_libraries( traccc_examples_cuda
    PUBLIC CUDA::cudart vecmem::core vecmem::cuda detray::core detray::detectors
-          traccc::core traccc::device_common traccc::cuda
+          traccc::core traccc::device_common traccc_cuda_utils traccc::cuda
           traccc_examples_common )
 
 traccc_add_executable( throughput_st_cuda "throughput_st.cpp"
    LINK_LIBRARIES indicators::indicators vecmem::core vecmem::cuda
                   detray::detectors detray::io traccc::io traccc::performance
-                  traccc::core traccc::device_common traccc::cuda
+                  traccc::core traccc::device_common traccc_cuda_utils traccc::cuda
                   traccc::options traccc_examples_cuda )
 
 traccc_add_executable( throughput_mt_cuda "throughput_mt.cpp"
    LINK_LIBRARIES indicators::indicators TBB::tbb vecmem::core vecmem::cuda
                   detray::detectors detray::io traccc::io traccc::performance
-                  traccc::core traccc::device_common traccc::cuda
+                  traccc::core traccc::device_common traccc_cuda_utils traccc::cuda
                   traccc::options traccc_examples_cuda )

extern/covfie/CMakeLists.txt

@@ -13,7 +13,7 @@ message( STATUS "Fetching covfie as part of the traccc project" )
 
 # Declare where to get covfie from.
 set( TRACCC_COVFIE_SOURCE
-   "URL;https://github.com/acts-project/covfie/archive/refs/tags/v0.15.2.tar.gz;URL_MD5;f17d9365abed550845e8b6b3708d39fb"
+   "URL;https://github.com/acts-project/covfie/archive/refs/tags/v0.15.4.tar.gz;URL_MD5;80de9b43644c59d8239fb8cc184c6276"
    CACHE STRING "Source for covfie, when built as part of this project" )
 mark_as_advanced( TRACCC_COVFIE_SOURCE )
 FetchContent_Declare( covfie SYSTEM ${TRACCC_COVFIE_SOURCE} )

tests/cuda/CMakeLists.txt

@@ -62,6 +62,7 @@ traccc_add_test(
     detray::test_common
     traccc::core
     traccc::device_common
+    traccc_cuda_utils
     traccc::cuda
     traccc::performance
     traccc::io