xref: /freebsd-src/contrib/llvm-project/clang/lib/Headers/__clang_hip_runtime_wrapper.h (revision 5f757f3ff9144b609b3c433dfd370cc6bdc191ad)
15ffd83dbSDimitry Andric /*===---- __clang_hip_runtime_wrapper.h - HIP runtime support ---------------===
25ffd83dbSDimitry Andric  *
35ffd83dbSDimitry Andric  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
45ffd83dbSDimitry Andric  * See https://llvm.org/LICENSE.txt for license information.
55ffd83dbSDimitry Andric  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
65ffd83dbSDimitry Andric  *
75ffd83dbSDimitry Andric  *===-----------------------------------------------------------------------===
85ffd83dbSDimitry Andric  */
95ffd83dbSDimitry Andric 
105ffd83dbSDimitry Andric /*
115ffd83dbSDimitry Andric  * WARNING: This header is intended to be directly -include'd by
125ffd83dbSDimitry Andric  * the compiler and is not supposed to be included by users.
135ffd83dbSDimitry Andric  *
145ffd83dbSDimitry Andric  */
155ffd83dbSDimitry Andric 
165ffd83dbSDimitry Andric #ifndef __CLANG_HIP_RUNTIME_WRAPPER_H__
175ffd83dbSDimitry Andric #define __CLANG_HIP_RUNTIME_WRAPPER_H__
185ffd83dbSDimitry Andric 
195ffd83dbSDimitry Andric #if __HIP__
205ffd83dbSDimitry Andric 
215ffd83dbSDimitry Andric #define __host__ __attribute__((host))
225ffd83dbSDimitry Andric #define __device__ __attribute__((device))
235ffd83dbSDimitry Andric #define __global__ __attribute__((global))
245ffd83dbSDimitry Andric #define __shared__ __attribute__((shared))
255ffd83dbSDimitry Andric #define __constant__ __attribute__((constant))
26fe6060f1SDimitry Andric #define __managed__ __attribute__((managed))
275ffd83dbSDimitry Andric 
28e8d8bef9SDimitry Andric #if !defined(__cplusplus) || __cplusplus < 201103L
29e8d8bef9SDimitry Andric   #define nullptr NULL;
30e8d8bef9SDimitry Andric #endif
31e8d8bef9SDimitry Andric 
32fe6060f1SDimitry Andric #ifdef __cplusplus
33fe6060f1SDimitry Andric extern "C" {
34fe6060f1SDimitry Andric   __attribute__((__visibility__("default")))
35fe6060f1SDimitry Andric   __attribute__((weak))
36fe6060f1SDimitry Andric   __attribute__((noreturn))
__cxa_pure_virtual(void)37fe6060f1SDimitry Andric   __device__ void __cxa_pure_virtual(void) {
38fe6060f1SDimitry Andric     __builtin_trap();
39fe6060f1SDimitry Andric   }
40fe6060f1SDimitry Andric   __attribute__((__visibility__("default")))
41fe6060f1SDimitry Andric   __attribute__((weak))
42fe6060f1SDimitry Andric   __attribute__((noreturn))
__cxa_deleted_virtual(void)43fe6060f1SDimitry Andric   __device__ void __cxa_deleted_virtual(void) {
44fe6060f1SDimitry Andric     __builtin_trap();
45fe6060f1SDimitry Andric   }
46fe6060f1SDimitry Andric }
47fe6060f1SDimitry Andric #endif //__cplusplus
48fe6060f1SDimitry Andric 
49fe6060f1SDimitry Andric #if !defined(__HIPCC_RTC__)
50*5f757f3fSDimitry Andric #if __has_include("hip/hip_version.h")
51*5f757f3fSDimitry Andric #include "hip/hip_version.h"
52*5f757f3fSDimitry Andric #endif // __has_include("hip/hip_version.h")
53*5f757f3fSDimitry Andric #endif // __HIPCC_RTC__
54*5f757f3fSDimitry Andric 
55*5f757f3fSDimitry Andric typedef __SIZE_TYPE__ __hip_size_t;
56*5f757f3fSDimitry Andric 
57*5f757f3fSDimitry Andric #ifdef __cplusplus
58*5f757f3fSDimitry Andric extern "C" {
59*5f757f3fSDimitry Andric #endif //__cplusplus
60*5f757f3fSDimitry Andric 
61*5f757f3fSDimitry Andric #if HIP_VERSION_MAJOR * 100 + HIP_VERSION_MINOR >= 405
62*5f757f3fSDimitry Andric __device__ unsigned long long __ockl_dm_alloc(unsigned long long __size);
63*5f757f3fSDimitry Andric __device__ void __ockl_dm_dealloc(unsigned long long __addr);
64*5f757f3fSDimitry Andric #if __has_feature(address_sanitizer)
65*5f757f3fSDimitry Andric __device__ unsigned long long __asan_malloc_impl(unsigned long long __size,
66*5f757f3fSDimitry Andric                                                  unsigned long long __pc);
67*5f757f3fSDimitry Andric __device__ void __asan_free_impl(unsigned long long __addr,
68*5f757f3fSDimitry Andric                                  unsigned long long __pc);
malloc(__hip_size_t __size)69*5f757f3fSDimitry Andric __attribute__((noinline, weak)) __device__ void *malloc(__hip_size_t __size) {
70*5f757f3fSDimitry Andric   unsigned long long __pc = (unsigned long long)__builtin_return_address(0);
71*5f757f3fSDimitry Andric   return (void *)__asan_malloc_impl(__size, __pc);
72*5f757f3fSDimitry Andric }
free(void * __ptr)73*5f757f3fSDimitry Andric __attribute__((noinline, weak)) __device__ void free(void *__ptr) {
74*5f757f3fSDimitry Andric   unsigned long long __pc = (unsigned long long)__builtin_return_address(0);
75*5f757f3fSDimitry Andric   __asan_free_impl((unsigned long long)__ptr, __pc);
76*5f757f3fSDimitry Andric }
77*5f757f3fSDimitry Andric #else // __has_feature(address_sanitizer)
malloc(__hip_size_t __size)78*5f757f3fSDimitry Andric __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {
79*5f757f3fSDimitry Andric   return (void *) __ockl_dm_alloc(__size);
80*5f757f3fSDimitry Andric }
free(void * __ptr)81*5f757f3fSDimitry Andric __attribute__((weak)) inline __device__ void free(void *__ptr) {
82*5f757f3fSDimitry Andric   __ockl_dm_dealloc((unsigned long long)__ptr);
83*5f757f3fSDimitry Andric }
84*5f757f3fSDimitry Andric #endif // __has_feature(address_sanitizer)
85*5f757f3fSDimitry Andric #else  // HIP version check
86*5f757f3fSDimitry Andric #if __HIP_ENABLE_DEVICE_MALLOC__
87*5f757f3fSDimitry Andric __device__ void *__hip_malloc(__hip_size_t __size);
88*5f757f3fSDimitry Andric __device__ void *__hip_free(void *__ptr);
89*5f757f3fSDimitry Andric __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {
90*5f757f3fSDimitry Andric   return __hip_malloc(__size);
91*5f757f3fSDimitry Andric }
92*5f757f3fSDimitry Andric __attribute__((weak)) inline __device__ void free(void *__ptr) {
93*5f757f3fSDimitry Andric   __hip_free(__ptr);
94*5f757f3fSDimitry Andric }
95*5f757f3fSDimitry Andric #else  // __HIP_ENABLE_DEVICE_MALLOC__
96*5f757f3fSDimitry Andric __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {
97*5f757f3fSDimitry Andric   __builtin_trap();
98*5f757f3fSDimitry Andric   return (void *)0;
99*5f757f3fSDimitry Andric }
100*5f757f3fSDimitry Andric __attribute__((weak)) inline __device__ void free(void *__ptr) {
101*5f757f3fSDimitry Andric   __builtin_trap();
102*5f757f3fSDimitry Andric }
103*5f757f3fSDimitry Andric #endif // __HIP_ENABLE_DEVICE_MALLOC__
104*5f757f3fSDimitry Andric #endif // HIP version check
105*5f757f3fSDimitry Andric 
106*5f757f3fSDimitry Andric #ifdef __cplusplus
107*5f757f3fSDimitry Andric } // extern "C"
108*5f757f3fSDimitry Andric #endif //__cplusplus
109*5f757f3fSDimitry Andric 
110*5f757f3fSDimitry Andric #if !defined(__HIPCC_RTC__)
111fe6060f1SDimitry Andric #include <cmath>
112fe6060f1SDimitry Andric #include <cstdlib>
113fe6060f1SDimitry Andric #include <stdlib.h>
11404eeddc0SDimitry Andric #if __has_include("hip/hip_version.h")
11504eeddc0SDimitry Andric #include "hip/hip_version.h"
11604eeddc0SDimitry Andric #endif // __has_include("hip/hip_version.h")
117fe6060f1SDimitry Andric #else
118fe6060f1SDimitry Andric typedef __SIZE_TYPE__ size_t;
119fe6060f1SDimitry Andric // Define macros which are needed to declare HIP device API's without standard
120fe6060f1SDimitry Andric // C/C++ headers. This is for readability so that these API's can be written
121fe6060f1SDimitry Andric // the same way as non-hipRTC use case. These macros need to be popped so that
122fe6060f1SDimitry Andric // they do not pollute users' name space.
123fe6060f1SDimitry Andric #pragma push_macro("NULL")
124fe6060f1SDimitry Andric #pragma push_macro("uint32_t")
125fe6060f1SDimitry Andric #pragma push_macro("uint64_t")
126fe6060f1SDimitry Andric #pragma push_macro("CHAR_BIT")
127fe6060f1SDimitry Andric #pragma push_macro("INT_MAX")
128fe6060f1SDimitry Andric #define NULL (void *)0
129fe6060f1SDimitry Andric #define uint32_t __UINT32_TYPE__
130fe6060f1SDimitry Andric #define uint64_t __UINT64_TYPE__
131fe6060f1SDimitry Andric #define CHAR_BIT __CHAR_BIT__
132fe6060f1SDimitry Andric #define INT_MAX __INTMAX_MAX__
133fe6060f1SDimitry Andric #endif // __HIPCC_RTC__
134fe6060f1SDimitry Andric 
1355ffd83dbSDimitry Andric #include <__clang_hip_libdevice_declares.h>
1365ffd83dbSDimitry Andric #include <__clang_hip_math.h>
137bdd1243dSDimitry Andric #include <__clang_hip_stdlib.h>
1385ffd83dbSDimitry Andric 
139fe6060f1SDimitry Andric #if defined(__HIPCC_RTC__)
140fe6060f1SDimitry Andric #include <__clang_hip_cmath.h>
141fe6060f1SDimitry Andric #else
1425ffd83dbSDimitry Andric #include <__clang_cuda_math_forward_declares.h>
143e8d8bef9SDimitry Andric #include <__clang_hip_cmath.h>
1445ffd83dbSDimitry Andric #include <__clang_cuda_complex_builtins.h>
1455ffd83dbSDimitry Andric #include <algorithm>
1465ffd83dbSDimitry Andric #include <complex>
1475ffd83dbSDimitry Andric #include <new>
148fe6060f1SDimitry Andric #endif // __HIPCC_RTC__
1495ffd83dbSDimitry Andric 
1505ffd83dbSDimitry Andric #define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1
151fe6060f1SDimitry Andric #if defined(__HIPCC_RTC__)
152fe6060f1SDimitry Andric #pragma pop_macro("NULL")
153fe6060f1SDimitry Andric #pragma pop_macro("uint32_t")
154fe6060f1SDimitry Andric #pragma pop_macro("uint64_t")
155fe6060f1SDimitry Andric #pragma pop_macro("CHAR_BIT")
156fe6060f1SDimitry Andric #pragma pop_macro("INT_MAX")
157fe6060f1SDimitry Andric #endif // __HIPCC_RTC__
1585ffd83dbSDimitry Andric #endif // __HIP__
1595ffd83dbSDimitry Andric #endif // __CLANG_HIP_RUNTIME_WRAPPER_H__
160