1 /*===---- __clang_hip_runtime_wrapper.h - HIP runtime support ---------------===
2  *
3  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4  * See https://llvm.org/LICENSE.txt for license information.
5  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6  *
7  *===-----------------------------------------------------------------------===
8  */
9 
10 /*
11  * WARNING: This header is intended to be directly -include'd by
12  * the compiler and is not supposed to be included by users.
13  *
14  */
15 
16 #ifndef __CLANG_HIP_RUNTIME_WRAPPER_H__
17 #define __CLANG_HIP_RUNTIME_WRAPPER_H__
18 
19 #if __HIP__
20 
21 #define __host__ __attribute__((host))
22 #define __device__ __attribute__((device))
23 #define __global__ __attribute__((global))
24 #define __shared__ __attribute__((shared))
25 #define __constant__ __attribute__((constant))
26 #define __managed__ __attribute__((managed))
27 
28 #if !defined(__cplusplus) || __cplusplus < 201103L
29   #define nullptr NULL;
30 #endif
31 
32 #ifdef __cplusplus
33 extern "C" {
34   __attribute__((__visibility__("default")))
35   __attribute__((weak))
36   __attribute__((noreturn))
37   __device__ void __cxa_pure_virtual(void) {
38     __builtin_trap();
39   }
40   __attribute__((__visibility__("default")))
41   __attribute__((weak))
42   __attribute__((noreturn))
43   __device__ void __cxa_deleted_virtual(void) {
44     __builtin_trap();
45   }
46 }
47 #endif //__cplusplus
48 
49 #if !defined(__HIPCC_RTC__)
50 #include <cmath>
51 #include <cstdlib>
52 #include <stdlib.h>
53 #if __has_include("hip/hip_version.h")
54 #include "hip/hip_version.h"
55 #endif // __has_include("hip/hip_version.h")
56 #else
57 typedef __SIZE_TYPE__ size_t;
58 // Define macros which are needed to declare HIP device API's without standard
59 // C/C++ headers. This is for readability so that these API's can be written
60 // the same way as non-hipRTC use case. These macros need to be popped so that
61 // they do not pollute users' name space.
62 #pragma push_macro("NULL")
63 #pragma push_macro("uint32_t")
64 #pragma push_macro("uint64_t")
65 #pragma push_macro("CHAR_BIT")
66 #pragma push_macro("INT_MAX")
67 #define NULL (void *)0
68 #define uint32_t __UINT32_TYPE__
69 #define uint64_t __UINT64_TYPE__
70 #define CHAR_BIT __CHAR_BIT__
71 #define INT_MAX __INTMAX_MAX__
72 #endif // __HIPCC_RTC__
73 
74 typedef __SIZE_TYPE__ __hip_size_t;
75 
76 #ifdef __cplusplus
77 extern "C" {
78 #endif //__cplusplus
79 
80 #if HIP_VERSION_MAJOR * 100 + HIP_VERSION_MINOR >= 405
81 extern "C" __device__ unsigned long long __ockl_dm_alloc(unsigned long long __size);
82 extern "C" __device__ void __ockl_dm_dealloc(unsigned long long __addr);
83 __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {
84   return (void *) __ockl_dm_alloc(__size);
85 }
86 __attribute__((weak)) inline __device__ void free(void *__ptr) {
87   __ockl_dm_dealloc((unsigned long long)__ptr);
88 }
89 #else  // HIP version check
90 #if __HIP_ENABLE_DEVICE_MALLOC__
91 __device__ void *__hip_malloc(__hip_size_t __size);
92 __device__ void *__hip_free(void *__ptr);
93 __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {
94   return __hip_malloc(__size);
95 }
96 __attribute__((weak)) inline __device__ void free(void *__ptr) {
97   __hip_free(__ptr);
98 }
99 #else
100 __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {
101   __builtin_trap();
102   return (void *)0;
103 }
104 __attribute__((weak)) inline __device__ void free(void *__ptr) {
105   __builtin_trap();
106 }
107 #endif
108 #endif // HIP version check
109 
110 #ifdef __cplusplus
111 } // extern "C"
112 #endif //__cplusplus
113 
114 #include <__clang_hip_libdevice_declares.h>
115 #include <__clang_hip_math.h>
116 
117 #if defined(__HIPCC_RTC__)
118 #include <__clang_hip_cmath.h>
119 #else
120 #include <__clang_cuda_math_forward_declares.h>
121 #include <__clang_hip_cmath.h>
122 #include <__clang_cuda_complex_builtins.h>
123 #include <algorithm>
124 #include <complex>
125 #include <new>
126 #endif // __HIPCC_RTC__
127 
128 #define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1
129 #if defined(__HIPCC_RTC__)
130 #pragma pop_macro("NULL")
131 #pragma pop_macro("uint32_t")
132 #pragma pop_macro("uint64_t")
133 #pragma pop_macro("CHAR_BIT")
134 #pragma pop_macro("INT_MAX")
135 #endif // __HIPCC_RTC__
136 #endif // __HIP__
137 #endif // __CLANG_HIP_RUNTIME_WRAPPER_H__
138