desul: Move lock_array_copied from global scope (#5999)

* desul: Move lock_array_copied fromg global scope * lock_array_copied->once
sandialabs · Mar 27, 2023 · 7c7ae9a · 7c7ae9a
1 parent a7a2d71
commit 7c7ae9a
Show file tree

Hide file tree

Showing 2 changed files with 8 additions and 16 deletions.
diff --git a/tpls/desul/include/desul/atomics/Lock_Array_CUDA.hpp b/tpls/desul/include/desul/atomics/Lock_Array_CUDA.hpp
@@ -108,27 +108,23 @@ __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
 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
@@ -115,27 +115,23 @@ __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
 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