Change all CUDA warp operations to synchronise all threads

alpaka-group · Dec 13, 2023 · 4b42a01 · 4b42a01
1 parent 6312741
commit 4b42a01
Showing 1 changed file with 8 additions and 8 deletions.
diff --git a/include/alpaka/warp/WarpUniformCudaHipBuiltIn.hpp b/include/alpaka/warp/WarpUniformCudaHipBuiltIn.hpp
@@ -1,4 +1,4 @@
-/* Copyright 2022 Sergei Bastrakov, David M. Rogers, Jan Stephan, Andrea Bocci, Bernhard Manfred Gruber, Aurora Perego
+/* Copyright 2023 Sergei Bastrakov, David M. Rogers, Jan Stephan, Andrea Bocci, Bernhard Manfred Gruber, Aurora Perego
  * SPDX-License-Identifier: MPL-2.0
  */
 
@@ -67,7 +67,7 @@ namespace alpaka::warp
                 std::int32_t predicate) -> std::int32_t
             {
 #        if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
-                return __all_sync(activemask(warp), predicate);
+                return __all_sync(0xffff'ffff, predicate);
 #        else
                 return __all(predicate);
 #        endif
@@ -82,7 +82,7 @@ namespace alpaka::warp
                 std::int32_t predicate) -> std::int32_t
             {
 #        if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
-                return __any_sync(activemask(warp), predicate);
+                return __any_sync(0xffff'ffff, predicate);
 #        else
                 return __any(predicate);
 #        endif
@@ -103,7 +103,7 @@ namespace alpaka::warp
 #        endif
             {
 #        if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
-                return __ballot_sync(activemask(warp), predicate);
+                return __ballot_sync(0xffff'ffff, predicate);
 #        else
                 return __ballot(predicate);
 #        endif
@@ -121,7 +121,7 @@ namespace alpaka::warp
                 std::int32_t width) -> T
             {
 #        if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
-                return __shfl_sync(activemask(warp), val, srcLane, width);
+                return __shfl_sync(0xffff'ffff, val, srcLane, width);
 #        else
                 return __shfl(val, srcLane, width);
 #        endif
@@ -139,7 +139,7 @@ namespace alpaka::warp
                 std::int32_t width) -> T
             {
 #        if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
-                return __shfl_up_sync(activemask(warp), val, offset, width);
+                return __shfl_up_sync(0xffff'ffff, val, offset, width);
 #        else
                 return __shfl_up(val, offset, width);
 #        endif
@@ -157,7 +157,7 @@ namespace alpaka::warp
                 std::int32_t width) -> T
             {
 #        if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
-                return __shfl_down_sync(activemask(warp), val, offset, width);
+                return __shfl_down_sync(0xffff'ffff, val, offset, width);
 #        else
                 return __shfl_down(val, offset, width);
 #        endif
@@ -175,7 +175,7 @@ namespace alpaka::warp
                 std::int32_t width) -> T
             {
 #        if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
-                return __shfl_xor_sync(activemask(warp), val, mask, width);
+                return __shfl_xor_sync(0xffff'ffff, val, mask, width);
 #        else
                 return __shfl_xor(val, mask, width);
 #        endif