diff --git a/tpls/desul/include/desul/atomics/Lock_Array_CUDA.hpp b/tpls/desul/include/desul/atomics/Lock_Array_CUDA.hpp index 4ff7196eed..b4dc4dae74 100644 --- a/tpls/desul/include/desul/atomics/Lock_Array_CUDA.hpp +++ b/tpls/desul/include/desul/atomics/Lock_Array_CUDA.hpp @@ -108,11 +108,6 @@ __device__ inline void unlock_address_cuda(void* ptr, desul::MemoryScopeNode) { atomicExch(&desul::Impl::CUDA_SPACE_ATOMIC_LOCKS_NODE[offset], 0); } -// Make lock_array_copied an explicit translation unit scope thingy -namespace { -static int lock_array_copied = 0; -} // namespace - #ifdef __CUDACC_RDC__ inline #else @@ -120,15 +115,16 @@ inline static #endif void copy_cuda_lock_arrays_to_device() { - if (lock_array_copied == 0) { + static bool once = []() { cudaMemcpyToSymbol(CUDA_SPACE_ATOMIC_LOCKS_DEVICE, &CUDA_SPACE_ATOMIC_LOCKS_DEVICE_h, sizeof(int32_t*)); cudaMemcpyToSymbol(CUDA_SPACE_ATOMIC_LOCKS_NODE, &CUDA_SPACE_ATOMIC_LOCKS_NODE_h, sizeof(int32_t*)); - } - lock_array_copied = 1; + return true; + }(); + (void)once; } } // namespace Impl diff --git a/tpls/desul/include/desul/atomics/Lock_Array_HIP.hpp b/tpls/desul/include/desul/atomics/Lock_Array_HIP.hpp index 53144fbc4c..b80e2d4599 100644 --- a/tpls/desul/include/desul/atomics/Lock_Array_HIP.hpp +++ b/tpls/desul/include/desul/atomics/Lock_Array_HIP.hpp @@ -115,11 +115,6 @@ __device__ inline void unlock_address_hip(void* ptr, desul::MemoryScopeNode) { atomicExch(&desul::Impl::HIP_SPACE_ATOMIC_LOCKS_NODE[offset], 0); } -// Make lock_array_copied an explicit translation unit scope thing -namespace { -static int lock_array_copied = 0; -} // namespace - #ifdef __CLANG_RDC__ inline #else @@ -127,15 +122,16 @@ inline static #endif void copy_hip_lock_arrays_to_device() { - if (lock_array_copied == 0) { + static bool once = []() { (void)hipMemcpyToSymbol(HIP_SYMBOL(HIP_SPACE_ATOMIC_LOCKS_DEVICE), &HIP_SPACE_ATOMIC_LOCKS_DEVICE_h, sizeof(int32_t*)); (void)hipMemcpyToSymbol(HIP_SYMBOL(HIP_SPACE_ATOMIC_LOCKS_NODE), &HIP_SPACE_ATOMIC_LOCKS_NODE_h, sizeof(int32_t*)); - } - lock_array_copied = 1; + return true; + }(); + (void)once; } } // namespace Impl