Make traccc::alpaka executables work in pure C++ (#832)

* Make traccc::alpaka executables work in pure C++: move device code to utils. Built on work by Ryan Cross

* Fix formatting

* Fix choice of CPU backend, move to alpaka device_copy in track_params_estimation

* Fix formatting

* Remove variable only used for one file

* Move to CDNA warp size for AMD

* Fix ifdef logic for SYCL

* Formatting

* Move back to vecmem::copy

* Move back to vecmem::copy

* Move tag logic to device_tag.hpp, make get_device_info return a string

* Remove position-independent executable logic: not needed now that this is pure C++

* Fix copy-paste error in alpaka cuda device_memory_resource

---------

Co-authored-by: Stewart Martin-Haugh stewart.martin-haugh@stfc.ac.uk <smh@cern.ch>

device/alpaka/CMakeLists.txt

@@ -27,6 +27,7 @@ traccc_add_alpaka_library( traccc_alpaka alpaka TYPE SHARED
   # Utility definitions.
   "include/traccc/alpaka/utils/make_prefix_sum_buff.hpp"
   "src/utils/make_prefix_sum_buff.cpp"
+  "src/utils/get_device_info.cpp"
   # Seed finding code.
   "include/traccc/alpaka/seeding/spacepoint_binning.hpp"
   "include/traccc/alpaka/seeding/seed_finding.hpp"

device/alpaka/include/traccc/alpaka/seeding/track_params_estimation.hpp

@@ -8,6 +8,8 @@
 #pragma once
 
 // Project include(s)
+#include <vecmem/utils/copy.hpp>
+
 #include "traccc/edm/seed.hpp"
 #include "traccc/edm/spacepoint.hpp"
 #include "traccc/edm/track_parameters.hpp"
@@ -58,7 +60,7 @@ struct track_params_estimation
     /// Memory resource used by the algorithm
     traccc::memory_resource m_mr;
     /// Copy object used by the algorithm
-    vecmem::copy& m_copy;
+    ::vecmem::copy& m_copy;
 };
 
 }  // namespace traccc::alpaka

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

@@ -0,0 +1,31 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2024 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+#include <alpaka/acc/Tag.hpp>
+
+namespace traccc::alpaka {
+
+// Get alpaka tag for current device
+#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
+using AccTag = ::alpaka::TagGpuCudaRt;
+#elif defined(ALPAKA_ACC_GPU_HIP_ENABLED)
+using AccTag = ::alpaka::TagGpuHipRt;
+#elif defined(ALPAKA_ACC_SYCL_ENABLED)
+#if defined(ALPAKA_SYCL_ONEAPI_CPU)
+using AccTag = ::alpaka::TagCpuSycl;
+#elif defined(ALPAKA_SYCL_ONEAPI_FPGA)
+using AccTag = ::alpaka::TagFpgaSyclIntel;
+#elif defined(ALPAKA_SYCL_ONEAPI_GPU)
+using AccTag = ::alpaka::TagGpuSyclIntel;
+#endif
+#elif defined(ALPAKA_ACC_CPU_B_SEQ_T_THREADS_ENABLED)
+using AccTag = ::alpaka::TagCpuThreads;
+#endif
+
+}  // namespace traccc::alpaka

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

@@ -0,0 +1,19 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2024 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+#include <string>
+
+namespace traccc::alpaka {
+
+/// Function that prints the current device information to the console.
+/// Included as part of the traccc::alpaka namespace, to avoid having to include
+/// alpaka headers in any users of the library.
+std::string get_device_info();
+
+}  // namespace traccc::alpaka

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

@@ -32,7 +32,7 @@
 #include <vecmem/utils/copy.hpp>
 #endif
 
-#include <alpaka/alpaka.hpp>
+#include "traccc/alpaka/utils/device_tag.hpp"
 
 // Forward declarations so we can compile the types below
 namespace vecmem {
@@ -69,7 +69,7 @@ struct host_device_types {
 };
 template <>
 struct host_device_types<::alpaka::TagGpuCudaRt> {
-    using device_memory_resource = ::vecmem::cuda::host_memory_resource;
+    using device_memory_resource = ::vecmem::cuda::device_memory_resource;
     using host_memory_resource = ::vecmem::cuda::host_memory_resource;
     using managed_memory_resource = ::vecmem::cuda::managed_memory_resource;
     using device_copy = ::vecmem::cuda::copy;
@@ -102,4 +102,13 @@ struct host_device_types<::alpaka::TagGpuSyclIntel> {
     using managed_memory_resource = ::vecmem::sycl::shared_memory_resource;
     using device_copy = ::vecmem::sycl::copy;
 };
+
+using device_memory_resource =
+    typename host_device_types<AccTag>::device_memory_resource;
+using host_memory_resource =
+    typename host_device_types<AccTag>::host_memory_resource;
+using managed_memory_resource =
+    typename host_device_types<AccTag>::managed_memory_resource;
+using device_copy = typename host_device_types<AccTag>::device_copy;
+
 }  // namespace traccc::alpaka::vecmem

device/alpaka/src/seeding/seed_finding.cpp

@@ -343,7 +343,9 @@ seed_finding::output_type seed_finding::operator()(
 
     // Calculate the number of threads and thread blocks to run the weight
     // updating kernel for.
-    threadsPerBlock = warpSize * 2 < maxThreads ? warpSize * 2 : maxThreads;
+    threadsPerBlock = getWarpSize<Acc>() * 2 < maxThreads
+                          ? getWarpSize<Acc>() * 2
+                          : maxThreads;
     blocksPerGrid =
         (pBufHost_counter->m_nTriplets + threadsPerBlock - 1) / threadsPerBlock;
     workDiv = makeWorkDiv<Acc>(blocksPerGrid, threadsPerBlock);

device/alpaka/src/seeding/track_params_estimation.cpp

@@ -63,7 +63,7 @@ track_params_estimation::output_type track_params_estimation::operator()(
     // Run the kernel
     ::alpaka::exec<Acc>(queue, workDiv, EstimateTrackParamsKernel{},
                         spacepoints_view, seeds_view, bfield, stddev,
-                        vecmem::get_data(params_buffer));
+                        ::vecmem::get_data(params_buffer));
     ::alpaka::wait(queue);
 
     return params_buffer;

device/alpaka/src/utils/get_device_info.cpp

@@ -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
+ */
+
+// Local include(s).
+#include "utils.hpp"
+
+// Project include(s).
+#include "traccc/alpaka/utils/get_device_info.hpp"
+
+namespace traccc::alpaka {
+
+std::string get_device_info() {
+    int device = 0;
+    auto devAcc = ::alpaka::getDevByIdx(::alpaka::Platform<Acc>{}, 0u);
+    return std::string("Using Alpaka device: " + ::alpaka::getName(devAcc) +
+                       " [id: " + std::to_string(device) + "] ");
+}
+
+}  // namespace traccc::alpaka

device/alpaka/src/utils/utils.hpp

@@ -8,34 +8,47 @@
 #pragma once
 
 #include <alpaka/alpaka.hpp>
-#include <alpaka/example/ExampleDefaultAcc.hpp>
-
-#ifdef ALPAKA_ACC_GPU_CUDA_ENABLED
-#include <vecmem/utils/cuda/copy.hpp>
-#endif
-
-#ifdef ALPAKA_ACC_GPU_HIP_ENABLED
-#include <vecmem/utils/hip/copy.hpp>
-#endif
-
-#include <vecmem/utils/copy.hpp>
 
 namespace traccc::alpaka {
 
 using Dim = ::alpaka::DimInt<1>;
 using Idx = uint32_t;
 using WorkDiv = ::alpaka::WorkDivMembers<Dim, Idx>;
 
-using Acc = ::alpaka::ExampleDefaultAcc<Dim, Idx>;
+// Get alpaka accelerator - based on alpaka/examples/ExampleDefaultAcc.hpp
+#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
+using Acc = ::alpaka::AccGpuCudaRt<Dim, Idx>;
+#elif defined(ALPAKA_ACC_GPU_HIP_ENABLED)
+using Acc = ::alpaka::AccGpuHipRt<Dim, Idx>;
+#elif defined(ALPAKA_ACC_SYCL_ENABLED)
+#if defined(ALPAKA_SYCL_ONEAPI_CPU)
+using Acc = ::alpaka::AccCpuSycl<Dim, Idx>;
+#elif defined(ALPAKA_SYCL_ONEAPI_FPGA)
+using Acc = ::alpaka::AccFpgaSyclIntel<Dim, Idx>;
+#elif defined(ALPAKA_SYCL_ONEAPI_GPU)
+using Acc = ::alpaka::AccGpuSyclIntel<Dim, Idx>;
+#endif
+#elif defined(ALPAKA_ACC_CPU_B_SEQ_T_THREADS_ENABLED)
+using Acc = ::alpaka::AccCpuThreads<Dim, Idx>;
+#else
+#error "No supported backend selected." //we definitely want to fail the build if no matching accelerator is found
+#endif
+
 using Host = ::alpaka::DevCpu;
 using Queue = ::alpaka::Queue<Acc, ::alpaka::Blocking>;
 
-static constexpr std::size_t warpSize =
-#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) || defined(ALPAKA_ACC_GPU_HIP_ENABLED)
-    32;
-#else
-    4;
-#endif
+template <typename TAcc>
+consteval std::size_t getWarpSize() {
+    if constexpr (::alpaka::accMatchesTags<TAcc, ::alpaka::TagGpuCudaRt,
+                                           ::alpaka::TagGpuSyclIntel>) {
+        return 32;
+    }
+    if constexpr (::alpaka::accMatchesTags<TAcc, ::alpaka::TagGpuHipRt>) {
+        return 64;
+    } else {
+        return 4;
+    }
+}
 
 template <typename TAcc>
 inline WorkDiv makeWorkDiv(Idx blocks, Idx threadsOrElements) {

examples/run/alpaka/CMakeLists.txt

@@ -15,16 +15,12 @@ include(traccc-alpaka-functions)
 traccc_enable_language_alpaka()
 
 if(alpaka_ACC_GPU_CUDA_ENABLE)
-  set_source_files_properties(${TRACCC_ALPAKA_EXAMPLE_SOURCES} PROPERTIES LANGUAGE CUDA)
-
   list (APPEND EXTRA_LIBS vecmem::cuda)
 elseif(alpaka_ACC_GPU_HIP_ENABLE)
   find_package( HIPToolkit REQUIRED )
-  set_source_files_properties(${TRACCC_ALPAKA_EXAMPLE_SOURCES} PROPERTIES LANGUAGE HIP)
   list(APPEND EXTRA_LIBS HIP::hiprt vecmem::hip)
 elseif(alpaka_ACC_SYCL_ENABLE)
   list(APPEND EXTRA_LIBS vecmem::sycl)
-  set_source_files_properties(${TRACCC_ALPAKA_EXAMPLE_SOURCES} PROPERTIES LANGUAGE SYCL)
 endif()
 
 set(LIBRARIES vecmem::core traccc::io traccc::performance
@@ -35,11 +31,3 @@ traccc_add_executable( seq_example_alpaka "seq_example_alpaka.cpp"
     LINK_LIBRARIES ${LIBRARIES} )
 traccc_add_executable( seeding_example_alpaka "seeding_example_alpaka.cpp"
     LINK_LIBRARIES ${LIBRARIES} )
-
-#Can only do this once target is defined, so need another if here
-if(alpaka_ACC_GPU_HIP_ENABLE)
-  set_target_properties( traccc_seq_example_alpaka PROPERTIES
-     POSITION_INDEPENDENT_CODE TRUE )
-  set_target_properties( traccc_seeding_example_alpaka PROPERTIES
-     POSITION_INDEPENDENT_CODE TRUE )
-endif()

examples/run/alpaka/seeding_example_alpaka.cpp

@@ -35,7 +35,6 @@
 #include "traccc/seeding/track_params_estimation.hpp"
 
 // Detray include(s).
-#include "alpaka/example/ExampleDefaultAcc.hpp"
 #include "detray/core/detector.hpp"
 #include "detray/detectors/bfield.hpp"
 #include "detray/io/frontend/detector_reader.hpp"
@@ -63,33 +62,19 @@ int seq_run(const traccc::opts::track_seeding& seeding_opts,
             const traccc::opts::performance& performance_opts,
             const traccc::opts::accelerator& accelerator_opts) {
 
-    using Dim = ::alpaka::DimInt<1>;
-    using Idx = uint32_t;
-
-    using Acc = ::alpaka::ExampleDefaultAcc<Dim, Idx>;
 #ifdef ALPAKA_ACC_SYCL_ENABLED
     ::sycl::queue q;
     vecmem::sycl::queue_wrapper qw{&q};
-    traccc::alpaka::vecmem::host_device_types<
-        ::alpaka::trait::AccToTag<Acc>::type>::device_copy copy(qw);
-    traccc::alpaka::vecmem::host_device_types<
-        ::alpaka::trait::AccToTag<Acc>::type>::host_memory_resource host_mr(qw);
-    traccc::alpaka::vecmem::host_device_types<
-        ::alpaka::trait::AccToTag<Acc>::type>::device_memory_resource
-        device_mr(qw);
-    traccc::alpaka::vecmem::host_device_types<
-        ::alpaka::trait::AccToTag<Acc>::type>::managed_memory_resource
-        mng_mr(qw);
+    traccc::alpaka::vecmem::device_copy copy(qw);
+    traccc::alpaka::vecmem::host_memory_resource host_mr(qw);
+    traccc::alpaka::vecmem::device_memory_resource device_mr(qw);
+    traccc::alpaka::vecmem::managed_memory_resource mng_mr(qw);
     traccc::memory_resource mr{device_mr, &host_mr};
 #else
-    traccc::alpaka::vecmem::host_device_types<
-        ::alpaka::trait::AccToTag<Acc>::type>::device_copy copy;
-    traccc::alpaka::vecmem::host_device_types<
-        ::alpaka::trait::AccToTag<Acc>::type>::host_memory_resource host_mr;
-    traccc::alpaka::vecmem::host_device_types<
-        ::alpaka::trait::AccToTag<Acc>::type>::device_memory_resource device_mr;
-    traccc::alpaka::vecmem::host_device_types<
-        ::alpaka::trait::AccToTag<Acc>::type>::managed_memory_resource mng_mr;
+    traccc::alpaka::vecmem::device_copy copy;
+    traccc::alpaka::vecmem::host_memory_resource host_mr;
+    traccc::alpaka::vecmem::device_memory_resource device_mr;
+    traccc::alpaka::vecmem::managed_memory_resource mng_mr;
     traccc::memory_resource mr{device_mr, &host_mr};
 #endif
 

examples/run/alpaka/seq_example_alpaka.cpp

@@ -6,7 +6,6 @@
  */
 
 // Project include(s).
-#include "alpaka/example/ExampleDefaultAcc.hpp"
 #include "traccc/alpaka/clusterization/clusterization_algorithm.hpp"
 #include "traccc/alpaka/clusterization/measurement_sorting_algorithm.hpp"
 #include "traccc/alpaka/seeding/seeding_algorithm.hpp"
@@ -66,28 +65,17 @@ int seq_run(const traccc::opts::detector& detector_opts,
     const traccc::vector3 field_vec = {0.f, 0.f,
                                        seeding_opts.seedfinder.bFieldInZ};
 
-    using Dim = ::alpaka::DimInt<1>;
-    using Idx = uint32_t;
-
-    using Acc = ::alpaka::ExampleDefaultAcc<Dim, Idx>;
     // Memory resources used by the application.
 #ifdef ALPAKA_ACC_SYCL_ENABLED
     ::sycl::queue q;
     vecmem::sycl::queue_wrapper qw{&q};
-    traccc::alpaka::vecmem::host_device_types<
-        alpaka::trait::AccToTag<Acc>::type>::device_copy copy(qw);
-    traccc::alpaka::vecmem::host_device_types<
-        alpaka::trait::AccToTag<Acc>::type>::host_memory_resource host_mr(qw);
-    traccc::alpaka::vecmem::host_device_types<
-        alpaka::trait::AccToTag<Acc>::type>::device_memory_resource
-        device_mr(qw);
+    traccc::alpaka::vecmem::device_copy copy(qw);
+    traccc::alpaka::vecmem::host_memory_resource host_mr(qw);
+    traccc::alpaka::vecmem::device_memory_resource device_mr(qw);
 #else
-    traccc::alpaka::vecmem::host_device_types<
-        alpaka::trait::AccToTag<Acc>::type>::device_copy copy;
-    traccc::alpaka::vecmem::host_device_types<
-        alpaka::trait::AccToTag<Acc>::type>::host_memory_resource host_mr;
-    traccc::alpaka::vecmem::host_device_types<
-        alpaka::trait::AccToTag<Acc>::type>::device_memory_resource device_mr;
+    traccc::alpaka::vecmem::device_copy copy;
+    traccc::alpaka::vecmem::host_memory_resource host_mr;
+    traccc::alpaka::vecmem::device_memory_resource device_mr;
 #endif
     traccc::memory_resource mr{device_mr, &host_mr};
 

tests/alpaka/CMakeLists.txt

@@ -4,28 +4,23 @@
 #
 # Mozilla Public License Version 2.0
 
-set(TRACCC_ALPAKA_TEST_SOURCES
-  alpaka_basic.cpp
-  test_cca.cpp
-)
-
-
 include(traccc-alpaka-functions)
 traccc_enable_language_alpaka()
 
 if(alpaka_ACC_GPU_CUDA_ENABLE)
-  set_source_files_properties(${TRACCC_ALPAKA_TEST_SOURCES} PROPERTIES LANGUAGE CUDA)
+  set_source_files_properties(alpaka_basic.cpp PROPERTIES LANGUAGE CUDA)
   list(APPEND DEVICE_LIBRARIES vecmem::cuda)
 elseif(alpaka_ACC_GPU_HIP_ENABLE)
-  set_source_files_properties(${TRACCC_ALPAKA_TEST_SOURCES} PROPERTIES LANGUAGE HIP)
+  set_source_files_properties(alpaka_basic.cpp PROPERTIES LANGUAGE HIP)
   list(APPEND DEVICE_LIBRARIES vecmem::hip)
 elseif(alpaka_ACC_SYCL_ENABLE)
   list(APPEND DEVICE_LIBRARIES vecmem::sycl)
-  set_source_files_properties(${TRACCC_ALPAKA_TEST_SOURCES} PROPERTIES LANGUAGE SYCL)
+  set_source_files_properties(alpaka_basic.cpp PROPERTIES LANGUAGE SYCL)
 endif()
 
 traccc_add_test( alpaka
-   ${TRACCC_ALPAKA_TEST_SOURCES}
+   alpaka_basic.cpp
+   test_cca.cpp
    LINK_LIBRARIES
    GTest::gtest_main
    traccc_tests_common

tests/alpaka/alpaka_basic.cpp

@@ -139,17 +139,13 @@ GTEST_TEST(AlpakaBasic, VecMemOp) {
 #ifdef ALPAKA_ACC_SYCL_ENABLED
     ::sycl::queue q;
     vecmem::sycl::queue_wrapper qw{&q};
-    traccc::alpaka::vecmem::host_device_types<
-        alpaka::trait::AccToTag<Acc>::type>::device_copy vm_copy(qw);
+    traccc::alpaka::vecmem::device_copy vm_copy(qw);
 #else
-    traccc::alpaka::vecmem::host_device_types<
-        alpaka::trait::AccToTag<Acc>::type>::device_copy vm_copy;
+    traccc::alpaka::vecmem::device_copy vm_copy;
 #endif
 
-    traccc::alpaka::vecmem::host_device_types<
-        alpaka::trait::AccToTag<Acc>::type>::host_memory_resource host_mr;
-    traccc::alpaka::vecmem::host_device_types<
-        alpaka::trait::AccToTag<Acc>::type>::device_memory_resource device_mr;
+    traccc::alpaka::vecmem::host_memory_resource host_mr;
+    traccc::alpaka::vecmem::device_memory_resource device_mr;
 
     vecmem::vector<float> host_vector{n, &host_mr};