avoid potential race condition HIP

nliber · Sep 11, 2023 · 96bb26b · 96bb26b
1 parent 9081d36
commit 96bb26b
Showing 1 changed file with 2 additions and 1 deletion.
diff --git a/core/src/HIP/Kokkos_HIP_ReduceScan.hpp b/core/src/HIP/Kokkos_HIP_ReduceScan.hpp
@@ -154,7 +154,8 @@ struct HIPReductionsFunctor<FunctorType, false> {
     int const lane_id =
         (threadIdx.y * blockDim.x + threadIdx.x) % HIPTraits::WarpSize;
     for (int delta = skip_vector ? blockDim.x : 1; delta < width; delta *= 2) {
-      if (lane_id + delta < HIPTraits::WarpSize) {
+      if (lane_id + delta < HIPTraits::WarpSize &&
+          (lane_id % (delta * 2) == 0)) {
         functor.join(value, value + delta);
       }
     }