Warning: This file is not a C or C++ file. It does not have highlighting.
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 | #if __has_include("hip/hip_version.h") |
51 | #include "hip/hip_version.h" |
52 | #endif // __has_include("hip/hip_version.h") |
53 | #endif // __HIPCC_RTC__ |
54 | |
55 | typedef __SIZE_TYPE__ __hip_size_t; |
56 | |
57 | #ifdef __cplusplus |
58 | extern "C" { |
59 | #endif //__cplusplus |
60 | |
61 | #if HIP_VERSION_MAJOR * 100 + HIP_VERSION_MINOR >= 405 |
62 | __device__ unsigned long long __ockl_dm_alloc(unsigned long long __size); |
63 | __device__ void __ockl_dm_dealloc(unsigned long long __addr); |
64 | #if __has_feature(address_sanitizer) |
65 | __device__ unsigned long long __asan_malloc_impl(unsigned long long __size, |
66 | unsigned long long __pc); |
67 | __device__ void __asan_free_impl(unsigned long long __addr, |
68 | unsigned long long __pc); |
69 | __attribute__((noinline, weak)) __device__ void *malloc(__hip_size_t __size) { |
70 | unsigned long long __pc = (unsigned long long)__builtin_return_address(0); |
71 | return (void *)__asan_malloc_impl(__size, __pc); |
72 | } |
73 | __attribute__((noinline, weak)) __device__ void free(void *__ptr) { |
74 | unsigned long long __pc = (unsigned long long)__builtin_return_address(0); |
75 | __asan_free_impl((unsigned long long)__ptr, __pc); |
76 | } |
77 | #else // __has_feature(address_sanitizer) |
78 | __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) { |
79 | return (void *) __ockl_dm_alloc(__size); |
80 | } |
81 | __attribute__((weak)) inline __device__ void free(void *__ptr) { |
82 | __ockl_dm_dealloc((unsigned long long)__ptr); |
83 | } |
84 | #endif // __has_feature(address_sanitizer) |
85 | #else // HIP version check |
86 | #if __HIP_ENABLE_DEVICE_MALLOC__ |
87 | __device__ void *__hip_malloc(__hip_size_t __size); |
88 | __device__ void *__hip_free(void *__ptr); |
89 | __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) { |
90 | return __hip_malloc(__size); |
91 | } |
92 | __attribute__((weak)) inline __device__ void free(void *__ptr) { |
93 | __hip_free(__ptr); |
94 | } |
95 | #else // __HIP_ENABLE_DEVICE_MALLOC__ |
96 | __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) { |
97 | __builtin_trap(); |
98 | return (void *)0; |
99 | } |
100 | __attribute__((weak)) inline __device__ void free(void *__ptr) { |
101 | __builtin_trap(); |
102 | } |
103 | #endif // __HIP_ENABLE_DEVICE_MALLOC__ |
104 | #endif // HIP version check |
105 | |
106 | #ifdef __cplusplus |
107 | } // extern "C" |
108 | #endif //__cplusplus |
109 | |
110 | #if !defined(__HIPCC_RTC__) |
111 | #include <cmath> |
112 | #include <cstdlib> |
113 | #include <stdlib.h> |
114 | #if __has_include("hip/hip_version.h") |
115 | #include "hip/hip_version.h" |
116 | #endif // __has_include("hip/hip_version.h") |
117 | #else |
118 | typedef __SIZE_TYPE__ size_t; |
119 | // Define macros which are needed to declare HIP device API's without standard |
120 | // C/C++ headers. This is for readability so that these API's can be written |
121 | // the same way as non-hipRTC use case. These macros need to be popped so that |
122 | // they do not pollute users' name space. |
123 | #pragma push_macro("NULL") |
124 | #pragma push_macro("uint32_t") |
125 | #pragma push_macro("uint64_t") |
126 | #pragma push_macro("CHAR_BIT") |
127 | #pragma push_macro("INT_MAX") |
128 | #define NULL (void *)0 |
129 | #define uint32_t __UINT32_TYPE__ |
130 | #define uint64_t __UINT64_TYPE__ |
131 | #define CHAR_BIT __CHAR_BIT__ |
132 | #define INT_MAX __INTMAX_MAX__ |
133 | #endif // __HIPCC_RTC__ |
134 | |
135 | #include <__clang_hip_libdevice_declares.h> |
136 | #include <__clang_hip_math.h> |
137 | #include <__clang_hip_stdlib.h> |
138 | |
139 | #if defined(__HIPCC_RTC__) |
140 | #include <__clang_hip_cmath.h> |
141 | #else |
142 | #include <__clang_cuda_math_forward_declares.h> |
143 | #include <__clang_hip_cmath.h> |
144 | #include <__clang_cuda_complex_builtins.h> |
145 | #include <algorithm> |
146 | #include <complex> |
147 | #include <new> |
148 | #endif // __HIPCC_RTC__ |
149 | |
150 | #define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1 |
151 | #if defined(__HIPCC_RTC__) |
152 | #pragma pop_macro("NULL") |
153 | #pragma pop_macro("uint32_t") |
154 | #pragma pop_macro("uint64_t") |
155 | #pragma pop_macro("CHAR_BIT") |
156 | #pragma pop_macro("INT_MAX") |
157 | #endif // __HIPCC_RTC__ |
158 | #endif // __HIP__ |
159 | #endif // __CLANG_HIP_RUNTIME_WRAPPER_H__ |
160 |
Warning: This file is not a C or C++ file. It does not have highlighting.