diff --git a/include/alpaka/warp/WarpUniformCudaHipBuiltIn.hpp b/include/alpaka/warp/WarpUniformCudaHipBuiltIn.hpp index 58b89d5f81d9..43106356c9f0 100644 --- 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(0xffffffff, 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(0xffffffff, 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(0xffffffff, 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(0xffffffff, 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(0xffffffff, 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(0xffffffff, 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(0xffffffff, val, mask, width); # else return __shfl_xor(val, mask, width); # endif