atomics.h

Go to the documentation of this file.

#pragma once
 
#include "backend/api.h"
 
namespace fz {
namespace backend {
 
#if defined(FZGMOD_BACKEND_HIP)
 
template <typename T>
__device__ inline T atomicAddBlock(T* addr, T val) {
    return __hip_atomic_fetch_add(addr, val, __ATOMIC_RELAXED,
                                  __HIP_MEMORY_SCOPE_WORKGROUP);
}
 
template <typename T>
__device__ inline T atomicOrBlock(T* addr, T val) {
    return __hip_atomic_fetch_or(addr, val, __ATOMIC_RELAXED,
                                 __HIP_MEMORY_SCOPE_WORKGROUP);
}
 
template <typename T>
__device__ inline T atomicMaxBlock(T* addr, T val) {
    return __hip_atomic_fetch_max(addr, val, __ATOMIC_RELAXED,
                                  __HIP_MEMORY_SCOPE_WORKGROUP);
}
 
#else // CUDA
 
template <typename T>
__device__ inline T atomicAddBlock(T* addr, T val) {
    return atomicAdd_block(addr, val);
}
 
template <typename T>
__device__ inline T atomicOrBlock(T* addr, T val) {
    return atomicOr_block(addr, val);
}
 
template <typename T>
__device__ inline T atomicMaxBlock(T* addr, T val) {
    return atomicMax_block(addr, val);
}
 
#endif
 
} // namespace backend
} // namespace fz