Merge pull request #1037 from stephenswat/perf/cuda_launch_config

Update CUDA propagation and fit kernel launch

Author: stephenswat

device/cuda/src/finding/combinatorial_kalman_filter.cuh

@@ -369,7 +369,7 @@ combinatorial_kalman_filter(
                     .tips_view = tips_buffer,
                     .tip_lengths_view = tip_length_buffer};
 
-                const unsigned int nThreads = warp_size * 2;
+                const unsigned int nThreads = warp_size * 4;
                 const unsigned int nBlocks =
                     (n_candidates + nThreads - 1) / nThreads;
                 propagate_to_next_surface<

device/cuda/src/finding/kernels/specializations/propagate_to_next_surface_src.cuh

@@ -18,7 +18,7 @@ namespace traccc::cuda {
 namespace kernels {
 
 template <typename propagator_t, typename bfield_t>
-__global__ void propagate_to_next_surface(
+__global__ __launch_bounds__(128, 8) void propagate_to_next_surface(
     const finding_config cfg,
     device::propagate_to_next_surface_payload<propagator_t, bfield_t> payload) {
 

device/cuda/src/fitting/kalman_fitting.cuh

@@ -109,7 +109,7 @@ track_state_container_types::buffer kalman_fitting(
     param_liveness_setup_event->ignore();
 
     // Launch parameters for all the kernels.
-    const unsigned int nThreads = warp_size * 2;
+    const unsigned int nThreads = warp_size * 4;
     const unsigned int nBlocks = (n_tracks + nThreads - 1) / nThreads;
 
     // Fill the keys and param_ids buffers.

device/cuda/src/fitting/kernels/specializations/fit_backward_src.cuh

@@ -14,8 +14,8 @@
 namespace traccc::cuda {
 namespace kernels {
 template <typename fitter_t>
-__global__ void fit_backward(const fitting_config cfg,
-                             const device::fit_payload<fitter_t> payload) {
+__global__ __launch_bounds__(128, 8) void fit_backward(
+    const fitting_config cfg, const device::fit_payload<fitter_t> payload) {
     device::fit_backward<fitter_t>(details::global_index1(), cfg, payload);
 }
 }  // namespace kernels

device/cuda/src/fitting/kernels/specializations/fit_forward_src.cuh

@@ -14,8 +14,8 @@
 namespace traccc::cuda {
 namespace kernels {
 template <typename fitter_t>
-__global__ void fit_forward(const fitting_config cfg,
-                            const device::fit_payload<fitter_t> payload) {
+__global__ __launch_bounds__(128, 8) void fit_forward(
+    const fitting_config cfg, const device::fit_payload<fitter_t> payload) {
     device::fit_forward<fitter_t>(details::global_index1(), cfg, payload);
 }
 }  // namespace kernels