Merge pull request #1283 from krasznaa/ClangFixes-main-20260330

Clang + FP64 Fixes, main branch (2026.03.30.)

Author: stephenswat

device/common/include/traccc/finding/device/impl/find_tracks.ipp

@@ -275,7 +275,8 @@ TRACCC_HOST_DEVICE inline void find_tracks(
                 const unsigned int owner_global_thread_id =
                     owner_local_thread_id +
                     thread_id.getBlockDimX() * thread_id.getBlockIdX();
-                const float chi2 = std::get<0>(*result).filtered_chi2();
+                const traccc::scalar chi2 =
+                    std::get<0>(*result).filtered_chi2();
                 assert(chi2 >= 0.f);
                 unsigned long long int* mutex_ptr =
                     &shared_payload
@@ -366,6 +367,8 @@ TRACCC_HOST_DEVICE inline void find_tracks(
                     const unsigned int p_offset =
                         owner_global_thread_id *
                         cfg.max_num_branches_per_surface;
+                    // No need to initialize this next variable. It always gets
+                    // a valid value in the proceeding expressions.
                     float new_max;
 
                     if (index == cfg.max_num_branches_per_surface) {
@@ -375,11 +378,13 @@ TRACCC_HOST_DEVICE inline void find_tracks(
                          * replace it. Also keep track of what the new maximum
                          * $\chi^2$ value will be.
                          */
-                        float highest = 0.f;
+                        traccc::scalar highest =
+                            static_cast<traccc::scalar>(0.f);
 
                         for (unsigned int i = 0;
                              i < cfg.max_num_branches_per_surface; ++i) {
-                            float old_chi2 = tmp_links.at(p_offset + i).chi2;
+                            const traccc::scalar old_chi2 =
+                                tmp_links.at(p_offset + i).chi2;
 
                             if (old_chi2 > highest) {
                                 highest = old_chi2;
@@ -390,14 +395,15 @@ TRACCC_HOST_DEVICE inline void find_tracks(
                         assert(l_pos !=
                                std::numeric_limits<unsigned int>::max());
 
-                        new_max = chi2;
+                        new_max = static_cast<float>(chi2);
 
                         for (unsigned int i = 0;
                              i < cfg.max_num_branches_per_surface; ++i) {
-                            float old_chi2 = tmp_links.at(p_offset + i).chi2;
+                            const traccc::scalar old_chi2 =
+                                tmp_links.at(p_offset + i).chi2;
 
                             if (i != l_pos && old_chi2 > new_max) {
-                                new_max = old_chi2;
+                                new_max = static_cast<float>(old_chi2);
                             }
 
                             assert(old_chi2 <=
@@ -407,7 +413,7 @@ TRACCC_HOST_DEVICE inline void find_tracks(
                         assert(chi2 <= new_max);
                     } else {
                         l_pos = index;
-                        new_max = std::max(chi2, max);
+                        new_max = std::max(static_cast<float>(chi2), max);
                     }
 
                     assert(l_pos < cfg.max_num_branches_per_surface);

device/common/include/traccc/finding/device/impl/gather_best_tips_per_measurement.ipp

@@ -121,12 +121,12 @@ TRACCC_HOST_DEVICE inline void gather_best_tips_per_measurement(
                     tip_index.at(offset + out_idx) = thread_id;
                     tip_pval.at(offset + out_idx) = pval;
 
-                    typename algebra_t::scalar new_worst =
-                        std::numeric_limits<typename algebra_t::scalar>::max();
+                    float new_worst = std::numeric_limits<float>::max();
 
                     for (unsigned int i = 0; i < new_size; ++i) {
-                        new_worst =
-                            std::min(new_worst, tip_pval.at(offset + i));
+                        new_worst = std::min(
+                            new_worst,
+                            static_cast<float>(tip_pval.at(offset + i)));
                     }
 
                     [[maybe_unused]] bool cas_result =

tests/cuda/test_grid2.cu

@@ -45,7 +45,7 @@ void grid_replace_test(grid2_view<host_grid2_replace> grid_view) {
     const auto& axis0 = grid_view._axis_p0_view;
     const auto& axis1 = grid_view._axis_p1_view;
 
-    int block_dim = 1;
+    unsigned int block_dim = 1;
     dim3 thread_dim(axis0.n_bins, axis1.n_bins);
 
     // run the kernel
@@ -83,7 +83,7 @@ void grid_replace_ci_test(grid2_view<host_grid2_replace_ci> grid_view) {
     const auto& axis0 = grid_view._axis_p0_view;
     const auto& axis1 = grid_view._axis_p1_view;
 
-    int block_dim = 1;
+    unsigned int block_dim = 1;
     dim3 thread_dim(axis0.n_bins, axis1.n_bins);
 
     // run the kernel
@@ -127,7 +127,7 @@ void grid_complete_test(grid2_view<host_grid2_complete> grid_view) {
     const auto& axis0 = grid_view._axis_p0_view;
     const auto& axis1 = grid_view._axis_p1_view;
 
-    int block_dim = 1;
+    unsigned int block_dim = 1;
     dim3 thread_dim(axis0.n_bins, axis1.n_bins);
 
     // run the kernel
@@ -162,7 +162,7 @@ void grid_attach_read_test(const_grid2_view<host_grid2_attach> grid_view) {
     const auto& axis0 = grid_view._axis_p0_view;
     const auto& axis1 = grid_view._axis_p1_view;
 
-    int block_dim = 1;
+    unsigned int block_dim = 1;
     dim3 thread_dim(axis0.n_bins, axis1.n_bins);
 
     // run the kernel
@@ -206,7 +206,7 @@ void grid_attach_fill_test(grid2_view<host_grid2_attach> grid_view) {
     const auto& axis1 = grid_view._axis_p1_view;
 
     dim3 block_dim(axis0.n_bins, axis1.n_bins);
-    int thread_dim = 100;
+    unsigned int thread_dim = 100;
 
     // run the kernel
     grid_attach_fill_test_kernel<<<block_dim, thread_dim>>>(grid_view);
@@ -237,7 +237,7 @@ __global__ void grid_attach_assign_test_kernel(
 // attach_fill_test implementation
 void grid_attach_assign_test(grid2_view<host_grid2_attach> grid_view) {
 
-    int block_dim = 1;
+    unsigned int block_dim = 1;
     dim3 thread_dim(2, 2);
 
     // run the kernel