1 //===---- hip_atomics.h - Declarations of hip atomic functions ---- C++ -*-===//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8
9 #ifndef OMPTARGET_AMDGCN_HIP_ATOMICS_H
10 #define OMPTARGET_AMDGCN_HIP_ATOMICS_H
11
12 #include "target_impl.h"
13
14 namespace {
15
atomicAdd(T * address,T val)16 template <typename T> DEVICE T atomicAdd(T *address, T val) {
17 return __atomic_fetch_add(address, val, __ATOMIC_SEQ_CST);
18 }
19
atomicMax(T * address,T val)20 template <typename T> DEVICE T atomicMax(T *address, T val) {
21 return __atomic_fetch_max(address, val, __ATOMIC_SEQ_CST);
22 }
23
atomicExch(T * address,T val)24 template <typename T> DEVICE T atomicExch(T *address, T val) {
25 T r;
26 __atomic_exchange(address, &val, &r, __ATOMIC_SEQ_CST);
27 return r;
28 }
29
atomicCAS(T * address,T compare,T val)30 template <typename T> DEVICE T atomicCAS(T *address, T compare, T val) {
31 (void)__atomic_compare_exchange(address, &compare, &val, false,
32 __ATOMIC_SEQ_CST, __ATOMIC_RELAXED);
33 return compare;
34 }
35
atomicInc(uint32_t * address,uint32_t max)36 INLINE uint32_t atomicInc(uint32_t *address, uint32_t max) {
37 return __builtin_amdgcn_atomic_inc32(address, max, __ATOMIC_SEQ_CST, "");
38 }
39
40 } // namespace
41 #endif
42