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 #if __has_feature(address_sanitizer) 84 extern "C" __device__ unsigned long long __asan_malloc_impl(unsigned long long __size, unsigned long long __pc); 85 extern "C" __device__ void __asan_free_impl(unsigned long long __addr, unsigned long long __pc); 86 __attribute__((noinline, weak)) __device__ void *malloc(__hip_size_t __size) { 87 unsigned long long __pc = (unsigned long long)__builtin_return_address(0); 88 return (void *)__asan_malloc_impl(__size, __pc); 89 } 90 __attribute__((noinline, weak)) __device__ void free(void *__ptr) { 91 unsigned long long __pc = (unsigned long long)__builtin_return_address(0); 92 __asan_free_impl((unsigned long long)__ptr, __pc); 93 } 94 #else 95 __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) { 96 return (void *) __ockl_dm_alloc(__size); 97 } 98 __attribute__((weak)) inline __device__ void free(void *__ptr) { 99 __ockl_dm_dealloc((unsigned long long)__ptr); 100 } 101 #endif // __has_feature(address_sanitizer) 102 #else // HIP version check 103 #if __HIP_ENABLE_DEVICE_MALLOC__ 104 __device__ void *__hip_malloc(__hip_size_t __size); 105 __device__ void *__hip_free(void *__ptr); 106 __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) { 107 return __hip_malloc(__size); 108 } 109 __attribute__((weak)) inline __device__ void free(void *__ptr) { 110 __hip_free(__ptr); 111 } 112 #else 113 __attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) { 114 __builtin_trap(); 115 return (void *)0; 116 } 117 __attribute__((weak)) inline __device__ void free(void *__ptr) { 118 __builtin_trap(); 119 } 120 #endif 121 #endif // HIP version check 122 123 #ifdef __cplusplus 124 } // extern "C" 125 #endif //__cplusplus 126 127 #include <__clang_hip_libdevice_declares.h> 128 #include <__clang_hip_math.h> 129 #include <__clang_hip_stdlib.h> 130 131 #if defined(__HIPCC_RTC__) 132 #include <__clang_hip_cmath.h> 133 #else 134 #include <__clang_cuda_math_forward_declares.h> 135 #include <__clang_hip_cmath.h> 136 #include <__clang_cuda_complex_builtins.h> 137 #include <algorithm> 138 #include <complex> 139 #include <new> 140 #endif // __HIPCC_RTC__ 141 142 #define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1 143 #if defined(__HIPCC_RTC__) 144 #pragma pop_macro("NULL") 145 #pragma pop_macro("uint32_t") 146 #pragma pop_macro("uint64_t") 147 #pragma pop_macro("CHAR_BIT") 148 #pragma pop_macro("INT_MAX") 149 #endif // __HIPCC_RTC__ 150 #endif // __HIP__ 151 #endif // __CLANG_HIP_RUNTIME_WRAPPER_H__ 152