text: /*
  //@HEADER
  // ************************************************************************
  //
  //                        Kokkos v. 2.0
  //              Copyright (2019) Sandia Corporation
  //
  // Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
  // the U.S. Government retains certain rights in this software.
  //
  // Redistribution and use in source and binary forms, with or without
  // modification, are permitted provided that the following conditions are
  // met:
  //
  // 1. Redistributions of source code must retain the above copyright
  // notice, this list of conditions and the following disclaimer.
  //
  // 2. Redistributions in binary form must reproduce the above copyright
  // notice, this list of conditions and the following disclaimer in the
  // documentation and/or other materials provided with the distribution.
  //
  // 3. Neither the name of the Corporation nor the names of the
  // contributors may be used to endorse or promote products derived from
  // this software without specific prior written permission.
  //
  // THIS SOFTWARE IS PROVIDED BY SANDIA CORPORATION "AS IS" AND ANY
  // EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
  // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
  // PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL SANDIA CORPORATION OR THE
  // CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
  // EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
  // PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
  // PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
  // LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
  // NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
  // SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
  //
  // Questions? Contact Christian R. Trott (crtrott@sandia.gov)
  //
  // ************************************************************************
  //@HEADER
  */
  
  #ifndef KOKKOS_IMPL_KOKKOS_ATOMIC_LOAD_HPP
  #define KOKKOS_IMPL_KOKKOS_ATOMIC_LOAD_HPP
  
  #include <Kokkos_Macros.hpp>
  #if defined(KOKKOS_ATOMIC_HPP)
  
  #include <impl/Kokkos_Atomic_Memory_Order.hpp>
  
  #if defined(KOKKOS_ENABLE_CUDA)
  #include <Cuda/Kokkos_Cuda_Atomic_Intrinsics.hpp>
  #endif
  
  namespace Kokkos {
  namespace Impl {
  
  // Olivier's implementation helpfully binds to the same builtins as GNU, so
  // we make this code common across multiple options
  #if (defined(KOKKOS_ENABLE_GNU_ATOMICS) && !defined(__CUDA_ARCH__)) \
      || (defined(KOKKOS_ENABLE_INTEL_ATOMICS) && !defined(__CUDA_ARCH__)) \
      || defined(KOKKOS_ENABLE_CUDA_ASM_ATOMICS)
  
  #if defined(__CUDA_ARCH__) && defined(KOKKOS_ENABLE_CUDA_ASM_ATOMICS)
    #define KOKKOS_INTERNAL_INLINE_DEVICE_IF_CUDA_ARCH __inline__ __device__
  #else
    #define KOKKOS_INTERNAL_INLINE_DEVICE_IF_CUDA_ARCH inline
  #endif
  
  template <class T, class MemoryOrder>
  KOKKOS_INTERNAL_INLINE_DEVICE_IF_CUDA_ARCH
  T _atomic_load(
    T* ptr, MemoryOrder,
    typename std::enable_if<
      (
        sizeof(T) == 1
          || sizeof(T) == 2
          || sizeof(T) == 4
          || sizeof(T) == 8
      )
      && std::is_same<
        typename MemoryOrder::memory_order,
        typename std::remove_cv<MemoryOrder>::type
      >::value,
      void const**
    >::type = nullptr
  ) {
    return __atomic_load_n(ptr, MemoryOrder::gnu_constant);
  }
  
  template <class T, class MemoryOrder>
  KOKKOS_INTERNAL_INLINE_DEVICE_IF_CUDA_ARCH
  T _atomic_load(
    T* ptr, MemoryOrder,
    typename std::enable_if<
      !(
        sizeof(T) == 1
          || sizeof(T) == 2
          || sizeof(T) == 4
          || sizeof(T) == 8
      )
      && std::is_default_constructible<T>::value
      && std::is_same<
        typename MemoryOrder::memory_order,
        typename std::remove_cv<MemoryOrder>::type
      >::value,
      void const**
    >::type = nullptr
  ) {
    T rv{};
    __atomic_load(ptr, &rv, MemoryOrder::gnu_constant);
    return rv;
  }
  
  #undef KOKKOS_INTERNAL_INLINE_DEVICE_IF_CUDA_ARCH
  
  #elif defined(__CUDA_ARCH__)
  
  // Not compiling for Volta or later, or Cuda ASM atomics were manually disabled
  
  template <class T>
  __device__ __inline__
  T _relaxed_atomic_load_impl(
    T* ptr,
    typename std::enable_if<
      (
        sizeof(T) == 1
        || sizeof(T) == 2
        || sizeof(T) == 4
        || sizeof(T) == 8
      ),
      void const**
    >::type = nullptr
  ) {
    return *ptr;
  }
  
  template <class T>
  struct NoOpOper {
    __device__ __inline__
    static constexpr T apply(T const&, T const&) noexcept { }
  };
  
  template <class T>
  __device__ __inline__
  T _relaxed_atomic_load_impl(
    T* ptr,
    typename std::enable_if<
      !(
        sizeof(T) == 1
        || sizeof(T) == 2
        || sizeof(T) == 4
        || sizeof(T) == 8
      ),
      void const**
    >::type = nullptr
  ) {
    T rv{};
    // TODO remove a copy operation here?
    Kokkos::atomic_oper_fetch(NoOpOper<T>{}, &rv, rv);
    return rv;
  }
  
  template <class T>
  __device__ __inline__
  T _atomic_load(T* ptr, memory_order_seq_cst_t) {
    Kokkos::memory_fence();
    T rv = Impl::_relaxed_atomic_load_impl(ptr);
    Kokkos::memory_fence();
    return rv;
  }
  
  template <class T>
  __device__ __inline__
  T _atomic_load(T* ptr, memory_order_acquire_t) {
    T rv = Impl::_relaxed_atomic_load_impl(ptr);
    Kokkos::memory_fence();
    return rv;
  }
  
  template <class T>
  __device__ __inline__
  T _atomic_load(T* ptr, memory_order_relaxed_t) {
    return _relaxed_atomic_load_impl(ptr);
  }
  
  #elif defined(KOKKOS_ENABLE_OPENMP_ATOMICS)
  
  template <class T, class MemoryOrder>
  inline
  T _atomic_load(T* ptr, MemoryOrder)
  {
    // AFAICT, all OpenMP atomics are sequentially consistent, so memory order doesn't matter
    T retval{ };
  #pragma omp atomic read
    {
      retval = *ptr;
    }
    return retval;
  }
  
  #elif defined(KOKKOS_ENABLE_SERIAL_ATOMICS)
  
  template <class T, class MemoryOrder>
  inline
  T _atomic_load(T* ptr, MemoryOrder)
  {
    return *ptr;
  }
  
  #endif // end of all atomic implementations
  
  
  template <class T>
  KOKKOS_FORCEINLINE_FUNCTION
  T atomic_load(T* ptr, Impl::memory_order_seq_cst_t) {
    return _atomic_load(ptr, Impl::memory_order_seq_cst);
  }
  
  template <class T>
  KOKKOS_FORCEINLINE_FUNCTION
  T atomic_load(T* ptr, Impl::memory_order_acquire_t) {
    return _atomic_load(ptr, Impl::memory_order_acquire);
  }
  
  template <class T>
  KOKKOS_FORCEINLINE_FUNCTION
  T atomic_load(T* ptr, Impl::memory_order_relaxed_t) {
    return _atomic_load(ptr, Impl::memory_order_relaxed);
  }
  
  template <class T>
  KOKKOS_FORCEINLINE_FUNCTION
  T atomic_load(T* ptr, Impl::memory_order_release_t) {
    static_assert(
      sizeof(T) == 0, // just something that will always be false, but only on instantiation
      "atomic_load with memory order release doesn't make any sense!"
    );
  }
  
  template <class T>
  KOKKOS_FORCEINLINE_FUNCTION
  T atomic_load(T* ptr, Impl::memory_order_acq_rel_t) {
    static_assert(
      sizeof(T) == 0, // just something that will always be false, but only on instantiation
      "atomic_load with memory order acq_rel doesn't make any sense!"
    );
  }
  
  template <class T>
  KOKKOS_FORCEINLINE_FUNCTION
  T atomic_load(T* ptr) {
    // relaxed by default!
    return _atomic_load(ptr, Impl::memory_order_relaxed);
  }
  
  } // end namespace Impl
  } // end namespace Kokkos
  
  #if defined(KOKKOS_ENABLE_CUDA)
  #include <Cuda/Kokkos_Cuda_Atomic_Intrinsics_Restore_Builtins.hpp>
  #endif
  
  #endif // defined(KOKKOS_ATOMIC_HPP)
  #endif //KOKKOS_IMPL_KOKKOS_ATOMIC_LOAD_HPP
  