1*bdd1243dSDimitry Andric /*===---- __clang_hip_stdlib.h - Device-side HIP math support --------------=== 2*bdd1243dSDimitry Andric * 3*bdd1243dSDimitry Andric * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4*bdd1243dSDimitry Andric * See https://llvm.org/LICENSE.txt for license information. 5*bdd1243dSDimitry Andric * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6*bdd1243dSDimitry Andric * 7*bdd1243dSDimitry Andric *===-----------------------------------------------------------------------=== 8*bdd1243dSDimitry Andric */ 9*bdd1243dSDimitry Andric #ifndef __CLANG_HIP_STDLIB_H__ 10*bdd1243dSDimitry Andric 11*bdd1243dSDimitry Andric #if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__) 12*bdd1243dSDimitry Andric #error "This file is for HIP and OpenMP AMDGCN device compilation only." 13*bdd1243dSDimitry Andric #endif 14*bdd1243dSDimitry Andric 15*bdd1243dSDimitry Andric #if !defined(__cplusplus) 16*bdd1243dSDimitry Andric 17*bdd1243dSDimitry Andric #include <limits.h> 18*bdd1243dSDimitry Andric 19*bdd1243dSDimitry Andric #ifdef __OPENMP_AMDGCN__ 20*bdd1243dSDimitry Andric #define __DEVICE__ static inline __attribute__((always_inline, nothrow)) 21*bdd1243dSDimitry Andric #else 22*bdd1243dSDimitry Andric #define __DEVICE__ static __device__ inline __attribute__((always_inline)) 23*bdd1243dSDimitry Andric #endif 24*bdd1243dSDimitry Andric 25*bdd1243dSDimitry Andric __DEVICE__ abs(int __x)26*bdd1243dSDimitry Andricint abs(int __x) { 27*bdd1243dSDimitry Andric int __sgn = __x >> (sizeof(int) * CHAR_BIT - 1); 28*bdd1243dSDimitry Andric return (__x ^ __sgn) - __sgn; 29*bdd1243dSDimitry Andric } 30*bdd1243dSDimitry Andric __DEVICE__ labs(long __x)31*bdd1243dSDimitry Andriclong labs(long __x) { 32*bdd1243dSDimitry Andric long __sgn = __x >> (sizeof(long) * CHAR_BIT - 1); 33*bdd1243dSDimitry Andric return (__x ^ __sgn) - __sgn; 34*bdd1243dSDimitry Andric } 35*bdd1243dSDimitry Andric __DEVICE__ llabs(long long __x)36*bdd1243dSDimitry Andriclong long llabs(long long __x) { 37*bdd1243dSDimitry Andric long long __sgn = __x >> (sizeof(long long) * CHAR_BIT - 1); 38*bdd1243dSDimitry Andric return (__x ^ __sgn) - __sgn; 39*bdd1243dSDimitry Andric } 40*bdd1243dSDimitry Andric 41*bdd1243dSDimitry Andric #endif // !defined(__cplusplus) 42*bdd1243dSDimitry Andric 43*bdd1243dSDimitry Andric #endif // #define __CLANG_HIP_STDLIB_H__ 44