Fix warp divergence (#681)

This PR fixes a crash observe while running benchmarks.



## Summary by CodeRabbit

* **Refactor**
  * Enhanced GPU kernel synchronization to improve performance and reliability in feasibility checking operations.

<sub>✏️ Tip: You can customize this high-level summary in your review settings.</sub>

Authors:
  - Hugo Linsenmaier (https://github.com/hlinsen)
  - Alice Boucher (https://github.com/aliceb-nv)

Approvers:
  - Chris Maes (https://github.com/chris-maes)
  - Alice Boucher (https://github.com/aliceb-nv)

URL: #681

Author: hlinsen

cpp/src/mip/feasibility_jump/feasibility_jump_kernels.cu

@@ -975,6 +975,7 @@ __device__ void compute_mtm_moves(typename fj_t<i_t, f_t>::climber_data_t::view_
   // related variable table couldn't be computed ahead of time, get related variables dynamically
   else if (fj.pb.related_variables.size() == 0) {
     compute_iteration_related_variables<i_t, f_t>(fj);
+    __syncwarp();
     cg::this_grid().sync();
     split_begin = 0;
     split_end   = fj.pb.n_variables;
@@ -1195,6 +1196,7 @@ DI thrust::tuple<i_t, f_t, typename fj_t<i_t, f_t>::move_score_t> gridwide_reduc
 
   // grid-wide reduce
   // will be replaced by a proper load balancing scheme
+  __syncwarp();
   cg::this_grid().sync();
 
   if (blockIdx.x == 0) {
@@ -1365,6 +1367,7 @@ __global__ void handle_local_minimum_kernel(typename fj_t<i_t, f_t>::climber_dat
 
   // Pick the best move among the variables involved in a random violated constraint.
   if (!fj.violated_constraints.empty()) {
+    __syncwarp();
     cg::this_grid().sync();
     thrust::tie(best_var, best_delta, best_score) =
       best_random_mtm_move<i_t, f_t, TPB_localmin>(fj);
@@ -1377,6 +1380,7 @@ __global__ void handle_local_minimum_kernel(typename fj_t<i_t, f_t>::climber_dat
   // also consider breakthrough moves
   if (*fj.best_objective < std::numeric_limits<f_t>::infinity() &&
       *fj.incumbent_objective > *fj.best_objective) {
+    __syncwarp();
     cg::this_grid().sync();
     auto [bm_best_var, bm_best_delta, bm_best_score] =
       best_breakthrough_move_at_local_min<i_t, f_t, TPB_localmin>(fj);
@@ -1392,6 +1396,7 @@ __global__ void handle_local_minimum_kernel(typename fj_t<i_t, f_t>::climber_dat
   }
 
   if (FIRST_THREAD) *fj.selected_var = best_var;
+  __syncwarp();
   cg::this_grid().sync();
   // still nothing? try sat MTM moves if we are in the feasible region
   // Attempt to find a valid move by going over MTM moves in valid constraints