The CUDA navigator test is failing in single precision, which has not been picked up by the CI. (acts-project#649)

Turns out the test does not configure the grids correctly, which can lead to the track hitting empty bins
even in straight line navigation.

Also sets the CI CUDA test to single precision.

Author: niermann999

.gitlab-ci.yml

@@ -16,7 +16,8 @@ build_cuda:
       - >
         cmake -S src -B build 
         -DCMAKE_BUILD_TYPE=Release
-        -DBUILD_TESTING=ON 
+        -DDETRAY_CUSTOM_SCALARTYPE=float
+        -DBUILD_TESTING=ON
         -DDETRAY_BUILD_TESTING=ON
         -DDETRAY_BUILD_CUDA=ON
         -DDETRAY_VC_PLUGIN=OFF
@@ -58,6 +59,7 @@ build_sycl:
       - >
         cmake -S src -B build 
         -DCMAKE_BUILD_TYPE=Release
+        -DDETRAY_CUSTOM_SCALARTYPE=float
         -DDETRAY_BUILD_CUDA=OFF
         -DDETRAY_BUILD_SYCL=ON
         -DBUILD_TESTING=ON 

tests/unit_tests/device/cuda/navigator_cuda.cpp

@@ -31,6 +31,8 @@ TEST(navigator_cuda, navigator) {
 
     // Create navigator
     navigator_host_t nav;
+    navigation::config cfg{};
+    cfg.search_window = {1u, 1u};
 
     // Create the vector of initial track parameters
     vecmem::vector<free_track_parameters<transform3>> tracks_host(&mng_mr);
@@ -60,21 +62,22 @@ TEST(navigator_cuda, navigator) {
         auto& track = tracks_host[i];
         stepper_t stepper;
 
+        // Propagator is built from the stepper and navigator
         prop_state<navigator_host_t::state> propagation{
             stepper_t::state{track}, navigator_host_t::state(det, mng_mr)};
 
         navigator_host_t::state& navigation = propagation._navigation;
         stepper_t::state& stepping = propagation._stepping;
 
         // Start propagation and record volume IDs
-        bool heartbeat = nav.init(propagation);
+        bool heartbeat = nav.init(propagation, cfg);
         while (heartbeat) {
 
             heartbeat &= stepper.step(propagation);
 
             navigation.set_high_trust();
 
-            heartbeat &= nav.update(propagation);
+            heartbeat &= nav.update(propagation, cfg);
 
             // Record volume
             volume_records_host[i].push_back(navigation.volume());
@@ -116,7 +119,7 @@ TEST(navigator_cuda, navigator) {
     copy.setup(candidates_buffer);
 
     // Run navigator test
-    navigator_test(det_data, tracks_data, candidates_buffer,
+    navigator_test(det_data, cfg, tracks_data, candidates_buffer,
                    volume_records_buffer, position_records_buffer);
 
     // Copy volume record buffer into volume & position records device

tests/unit_tests/device/cuda/navigator_cuda_kernel.cu

@@ -11,7 +11,7 @@
 namespace detray {
 
 __global__ void navigator_test_kernel(
-    typename detector_host_t::view_type det_data,
+    typename detector_host_t::view_type det_data, navigation::config cfg,
     vecmem::data::vector_view<free_track_parameters<transform3>> tracks_data,
     vecmem::data::jagged_vector_view<intersection_t> candidates_data,
     vecmem::data::jagged_vector_view<dindex> volume_records_data,
@@ -47,14 +47,14 @@ __global__ void navigator_test_kernel(
     navigation.set_volume(0u);
 
     // Start propagation and record volume IDs
-    bool heartbeat = nav.init(propagation);
+    bool heartbeat = nav.init(propagation, cfg);
     while (heartbeat) {
 
         heartbeat &= stepper.step(propagation);
 
         navigation.set_high_trust();
 
-        heartbeat = nav.update(propagation);
+        heartbeat = nav.update(propagation, cfg);
 
         // Record volume
         volume_records[gid].push_back(navigation.volume());
@@ -63,7 +63,7 @@ __global__ void navigator_test_kernel(
 }
 
 void navigator_test(
-    typename detector_host_t::view_type det_data,
+    typename detector_host_t::view_type det_data, navigation::config& cfg,
     vecmem::data::vector_view<free_track_parameters<transform3>>& tracks_data,
     vecmem::data::jagged_vector_view<intersection_t>& candidates_data,
     vecmem::data::jagged_vector_view<dindex>& volume_records_data,
@@ -74,7 +74,7 @@ void navigator_test(
 
     // run the test kernel
     navigator_test_kernel<<<block_dim, thread_dim>>>(
-        det_data, tracks_data, candidates_data, volume_records_data,
+        det_data, cfg, tracks_data, candidates_data, volume_records_data,
         position_records_data);
 
     // cuda error check

tests/unit_tests/device/cuda/navigator_cuda_kernel.hpp

@@ -40,23 +40,19 @@ constexpr unsigned int theta_steps{100u};
 constexpr unsigned int phi_steps{100u};
 
 constexpr scalar pos_diff_tolerance{1e-3f};
-constexpr scalar overstep_tolerance{-1e-4f};
 
 // dummy propagator state
 template <typename navigation_t>
 struct prop_state {
     stepper_t::state _stepping;
     navigation_t _navigation;
-
-    DETRAY_HOST_DEVICE
-    scalar mask_tolerance() const { return 15.f * unit<scalar>::um; }
 };
 
 namespace detray {
 
 /// test function for navigator with single state
 void navigator_test(
-    typename detector_host_t::view_type det_data,
+    typename detector_host_t::view_type det_data, navigation::config& cfg,
     vecmem::data::vector_view<free_track_parameters<transform3>>& tracks_data,
     vecmem::data::jagged_vector_view<intersection_t>& candidates_data,
     vecmem::data::jagged_vector_view<dindex>& volume_records_data,