CUDA updates for the common track finding changes.

acts-project · Jan 7, 2025 · febca64 · febca64
1 parent 63dd4e5
commit febca64
Show file tree

Hide file tree

Showing 3 changed files with 8 additions and 11 deletions.
diff --git a/device/cuda/src/finding/finding_algorithm.cu b/device/cuda/src/finding/finding_algorithm.cu
@@ -1,6 +1,6 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2023-2024 CERN for the benefit of the ACTS project
+ * (c) 2023-2025 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
@@ -189,8 +189,8 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
 
             kernels::apply_interaction<std::decay_t<detector_type>>
                 <<<nBlocks, nThreads, 0, stream>>>(
-                    m_cfg, {det_view, static_cast<int>(n_in_params),
-                            in_params_buffer, param_liveness_buffer});
+                    m_cfg, {det_view, n_in_params, in_params_buffer,
+                            param_liveness_buffer});
             TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
         }
 

diff --git a/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh b/device/cuda/src/finding/kernels/specializations/find_tracks_src.cuh
@@ -1,6 +1,6 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2023-2024 CERN for the benefit of the ACTS project
+ * (c) 2023-2025 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
@@ -29,8 +29,7 @@ __global__ void find_tracks(const finding_config cfg,
     cuda::barrier barrier;
     cuda::thread_id1 thread_id;
 
-    device::find_tracks<cuda::thread_id1, cuda::barrier, detector_t,
-                        finding_config>(
+    device::find_tracks<detector_t>(
         thread_id, barrier, cfg, payload,
         {shared_num_candidates, shared_candidates, shared_candidates_size});
 }

diff --git a/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh b/device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh
@@ -1,6 +1,6 @@
 /** TRACCC library, part of the ACTS project (R&D line)
  *
- * (c) 2023-2024 CERN for the benefit of the ACTS project
+ * (c) 2023-2025 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
@@ -18,10 +18,8 @@ __global__ void propagate_to_next_surface(
     const finding_config cfg,
     device::propagate_to_next_surface_payload<propagator_t, bfield_t> payload) {
 
-    int gid = threadIdx.x + blockIdx.x * blockDim.x;
-
-    device::propagate_to_next_surface<propagator_t, bfield_t, finding_config>(
-        gid, cfg, payload);
+    device::propagate_to_next_surface<propagator_t, bfield_t>(
+        threadIdx.x + blockIdx.x * blockDim.x, cfg, payload);
 }
 
 }  // namespace traccc::cuda::kernels