6024e5c395
clang 20.0.0 (based on r547379) from build 12806354. Bug: http://b/379133546 Test: N/A Change-Id: I2eb8938af55d809de674be63cb30cf27e801862b Upstream-Commit: ad834e67b1105d15ef907f6255d4c96e8e733f57
44 lines
1.2 KiB
C
44 lines
1.2 KiB
C
/*===---- __clang_hip_stdlib.h - Device-side HIP math support --------------===
|
|
*
|
|
* Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
|
|
* See https://llvm.org/LICENSE.txt for license information.
|
|
* SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
|
|
*
|
|
*===-----------------------------------------------------------------------===
|
|
*/
|
|
#ifndef __CLANG_HIP_STDLIB_H__
|
|
|
|
#if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)
|
|
#error "This file is for HIP and OpenMP AMDGCN device compilation only."
|
|
#endif
|
|
|
|
#if !defined(__cplusplus)
|
|
|
|
#include <limits.h>
|
|
|
|
#ifdef __OPENMP_AMDGCN__
|
|
#define __DEVICE__ static inline __attribute__((always_inline, nothrow))
|
|
#else
|
|
#define __DEVICE__ static __device__ inline __attribute__((always_inline))
|
|
#endif
|
|
|
|
__DEVICE__
|
|
int abs(int __x) {
|
|
int __sgn = __x >> (sizeof(int) * CHAR_BIT - 1);
|
|
return (__x ^ __sgn) - __sgn;
|
|
}
|
|
__DEVICE__
|
|
long labs(long __x) {
|
|
long __sgn = __x >> (sizeof(long) * CHAR_BIT - 1);
|
|
return (__x ^ __sgn) - __sgn;
|
|
}
|
|
__DEVICE__
|
|
long long llabs(long long __x) {
|
|
long long __sgn = __x >> (sizeof(long long) * CHAR_BIT - 1);
|
|
return (__x ^ __sgn) - __sgn;
|
|
}
|
|
|
|
#endif // !defined(__cplusplus)
|
|
|
|
#endif // #define __CLANG_HIP_STDLIB_H__
|