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