diff --git a/tpls/desul/include/desul/atomics/Lock_Array_CUDA.hpp b/tpls/desul/include/desul/atomics/Lock_Array_CUDA.hpp index e514061ed0..4ff7196eed 100644 --- a/tpls/desul/include/desul/atomics/Lock_Array_CUDA.hpp +++ b/tpls/desul/include/desul/atomics/Lock_Array_CUDA.hpp @@ -9,13 +9,11 @@ SPDX-License-Identifier: (BSD-3-Clause) #ifndef DESUL_ATOMICS_LOCK_ARRAY_CUDA_HPP_ #define DESUL_ATOMICS_LOCK_ARRAY_CUDA_HPP_ +#include + #include "desul/atomics/Common.hpp" #include "desul/atomics/Macros.hpp" -#ifdef DESUL_HAVE_CUDA_ATOMICS - -#include - namespace desul { namespace Impl { @@ -42,14 +40,6 @@ void init_lock_arrays_cuda(); template void finalize_lock_arrays_cuda(); -} // namespace Impl -} // namespace desul - -#if defined(__CUDACC__) - -namespace desul { -namespace Impl { - /// \brief This global variable in CUDA space is what kernels use /// to get access to the lock arrays. /// @@ -118,12 +108,7 @@ __device__ inline void unlock_address_cuda(void* ptr, desul::MemoryScopeNode) { atomicExch(&desul::Impl::CUDA_SPACE_ATOMIC_LOCKS_NODE[offset], 0); } -} // namespace Impl -} // namespace desul - // Make lock_array_copied an explicit translation unit scope thingy -namespace desul { -namespace Impl { namespace { static int lock_array_copied = 0; } // namespace @@ -149,13 +134,9 @@ inline static } // namespace Impl } // namespace desul -#endif /* defined( __CUDACC__ ) */ - -#endif /* defined( DESUL_HAVE_CUDA_ATOMICS ) */ - namespace desul { -#if defined(__CUDACC_RDC__) || (!defined(__CUDACC__)) +#if defined(__CUDACC_RDC__) inline void ensure_cuda_lock_arrays_on_device() {} #else static inline void ensure_cuda_lock_arrays_on_device() { diff --git a/tpls/desul/include/desul/atomics/Lock_Array_HIP.hpp b/tpls/desul/include/desul/atomics/Lock_Array_HIP.hpp index 33450e32ec..9290aea2b3 100644 --- a/tpls/desul/include/desul/atomics/Lock_Array_HIP.hpp +++ b/tpls/desul/include/desul/atomics/Lock_Array_HIP.hpp @@ -43,12 +43,6 @@ void init_lock_arrays_hip(); /// snapshotted version while also linking against pure Desul template void finalize_lock_arrays_hip(); -} // namespace Impl -} // namespace desul - -#ifdef __HIPCC__ -namespace desul { -namespace Impl { /** * \brief This global variable in HIP space is what kernels use to get access @@ -120,13 +114,8 @@ __device__ inline void unlock_address_hip(void* ptr, desul::MemoryScopeNode) { offset = offset & HIP_SPACE_ATOMIC_MASK; atomicExch(&desul::Impl::HIP_SPACE_ATOMIC_LOCKS_NODE[offset], 0); } -#endif -} // namespace Impl -} // namespace desul // Make lock_array_copied an explicit translation unit scope thing -namespace desul { -namespace Impl { namespace { static int lock_array_copied = 0; } // namespace @@ -150,7 +139,7 @@ inline static } } // namespace Impl -#if defined(DESUL_HIP_RDC) || (!defined(__HIPCC__)) +#if defined(DESUL_HIP_RDC) inline void ensure_hip_lock_arrays_on_device() {} #else static inline void ensure_hip_lock_arrays_on_device() { diff --git a/tpls/desul/src/Lock_Array_CUDA.cpp b/tpls/desul/src/Lock_Array_CUDA.cpp index 19944b378e..d8ab895b2b 100644 --- a/tpls/desul/src/Lock_Array_CUDA.cpp +++ b/tpls/desul/src/Lock_Array_CUDA.cpp @@ -11,7 +11,6 @@ SPDX-License-Identifier: (BSD-3-Clause) #include #include -#ifdef DESUL_HAVE_CUDA_ATOMICS #ifdef __CUDACC_RDC__ namespace desul { namespace Impl { @@ -96,4 +95,3 @@ template void finalize_lock_arrays_cuda(); } // namespace Impl } // namespace desul -#endif diff --git a/tpls/desul/src/Lock_Array_HIP.cpp b/tpls/desul/src/Lock_Array_HIP.cpp index 986f5475ae..6191fe81e2 100644 --- a/tpls/desul/src/Lock_Array_HIP.cpp +++ b/tpls/desul/src/Lock_Array_HIP.cpp @@ -11,7 +11,6 @@ SPDX-License-Identifier: (BSD-3-Clause) #include #include -#ifdef DESUL_HAVE_HIP_ATOMICS #ifdef DESUL_HIP_RDC namespace desul { namespace Impl { @@ -99,4 +98,3 @@ template void finalize_lock_arrays_hip(); } // namespace Impl } // namespace desul -#endif