Adjust propagation kernel launch bounds (#1131)

This commit modifies the launch bounds for our propagation and fitting
kernels, removing the minimum number of blocks per SM. This helps
performance as it allows the kernel to use more registers. The
parallelism on these kernels is currently bound by local memory, anyway.

Author: stephenswat

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__ __launch_bounds__(128, 8) void propagate_to_next_surface(
+__global__ __launch_bounds__(128) void propagate_to_next_surface(
     const finding_config cfg,
     device::propagate_to_next_surface_payload<propagator_t, bfield_t> payload) {
 

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

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

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

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