| //===---- hip_atomics.h - Declarations of hip atomic functions ---- C++ -*-===// |
| // |
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. |
| // See https://llvm.org/LICENSE.txt for license information. |
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception |
| // |
| //===----------------------------------------------------------------------===// |
| |
| #ifndef OMPTARGET_AMDGCN_HIP_ATOMICS_H |
| #define OMPTARGET_AMDGCN_HIP_ATOMICS_H |
| |
| #include "target_impl.h" |
| |
| namespace { |
| |
| template <typename T> DEVICE T atomicAdd(T *address, T val) { |
| return __atomic_fetch_add(address, val, __ATOMIC_SEQ_CST); |
| } |
| |
| template <typename T> DEVICE T atomicMax(T *address, T val) { |
| return __atomic_fetch_max(address, val, __ATOMIC_SEQ_CST); |
| } |
| |
| template <typename T> DEVICE T atomicExch(T *address, T val) { |
| T r; |
| __atomic_exchange(address, &val, &r, __ATOMIC_SEQ_CST); |
| return r; |
| } |
| |
| template <typename T> DEVICE T atomicCAS(T *address, T compare, T val) { |
| (void)__atomic_compare_exchange(address, &compare, &val, false, |
| __ATOMIC_SEQ_CST, __ATOMIC_RELAXED); |
| return compare; |
| } |
| |
| INLINE uint32_t atomicInc(uint32_t *address, uint32_t max) { |
| return __builtin_amdgcn_atomic_inc32(address, max, __ATOMIC_SEQ_CST, ""); |
| } |
| |
| } // namespace |
| #endif |
