aboutsummaryrefslogtreecommitdiff
path: root/contrib/llvm-project/clang/lib/Headers/__clang_hip_stdlib.h
diff options
context:
space:
mode:
Diffstat (limited to 'contrib/llvm-project/clang/lib/Headers/__clang_hip_stdlib.h')
-rw-r--r--contrib/llvm-project/clang/lib/Headers/__clang_hip_stdlib.h43
1 files changed, 43 insertions, 0 deletions
diff --git a/contrib/llvm-project/clang/lib/Headers/__clang_hip_stdlib.h b/contrib/llvm-project/clang/lib/Headers/__clang_hip_stdlib.h
new file mode 100644
index 000000000000..bd770e2415f9
--- /dev/null
+++ b/contrib/llvm-project/clang/lib/Headers/__clang_hip_stdlib.h
@@ -0,0 +1,43 @@
+/*===---- __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__