Alpaka inhomogeneous magnetic field

.github/workflows/builds.yml

@@ -75,7 +75,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:
@@ -87,7 +87,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.

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 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,50 @@
 #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>>>>>>;
+#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,
+                                   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);
+#endif
+#else
+    return bfield;
+}
+
+}  // namespace traccc::alpaka

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>
@@ -32,7 +35,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>(

extern/covfie/CMakeLists.txt

@@ -13,10 +13,14 @@ 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"
+   "GIT_REPOSITORY;https://github.com/StewMH/covfie/archive/refs/tags/v0.15.2.tar.gz;URL_MD5;f17d9365abed550845e8b6b3708d39fb"
    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} )
+#FetchContent_Declare( covfie SYSTEM ${TRACCC_COVFIE_SOURCE} )
+FetchContent_Declare( covfie SYSTEM
+GIT_REPOSITORY https://github.com/StewMH/covfie.git
+  GIT_TAG      hip_target_features_fix
+)
 
 # Options used for covfie.
 set( COVFIE_BUILD_EXAMPLES OFF CACHE BOOL "Build covfie examples")