1303241Sdim/*===---- __clang_hip_stdlib.h - Device-side HIP math support --------------===
2303241Sdim *
3353358Sdim * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4353358Sdim * See https://llvm.org/LICENSE.txt for license information.
5353358Sdim * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6303241Sdim *
7303241Sdim *===-----------------------------------------------------------------------===
8303241Sdim */
9303241Sdim#ifndef __CLANG_HIP_STDLIB_H__
10303241Sdim
11303241Sdim#if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)
12303241Sdim#error "This file is for HIP and OpenMP AMDGCN device compilation only."
13303241Sdim#endif
14303241Sdim
15341825Sdim#if !defined(__cplusplus)
16341825Sdim
17314564Sdim#include <limits.h>
18303241Sdim
19303241Sdim#ifdef __OPENMP_AMDGCN__
20314564Sdim#define __DEVICE__ static inline __attribute__((always_inline, nothrow))
21303241Sdim#else
22314564Sdim#define __DEVICE__ static __device__ inline __attribute__((always_inline))
23303241Sdim#endif
24314564Sdim
25314564Sdim__DEVICE__
26314564Sdimint abs(int __x) {
27314564Sdim  int __sgn = __x >> (sizeof(int) * CHAR_BIT - 1);
28314564Sdim  return (__x ^ __sgn) - __sgn;
29314564Sdim}
30314564Sdim__DEVICE__
31314564Sdimlong labs(long __x) {
32314564Sdim  long __sgn = __x >> (sizeof(long) * CHAR_BIT - 1);
33314564Sdim  return (__x ^ __sgn) - __sgn;
34314564Sdim}
35314564Sdim__DEVICE__
36314564Sdimlong long llabs(long long __x) {
37314564Sdim  long long __sgn = __x >> (sizeof(long long) * CHAR_BIT - 1);
38314564Sdim  return (__x ^ __sgn) - __sgn;
39360784Sdim}
40314564Sdim
41360784Sdim#endif // !defined(__cplusplus)
42360784Sdim
43360784Sdim#endif // #define __CLANG_HIP_STDLIB_H__
44314564Sdim