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)) __cxa_pure_virtual(void)37 __device__ void __cxa_pure_virtual(void) { 38 __builtin_trap(); 39 } 40 __attribute__((__visibility__("default"))) 41 __attribute__((weak)) 42 __attribute__((noreturn)) __cxa_deleted_virtual(void)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); malloc(__hip_size_t __size)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 } free(void * __ptr)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) malloc(__hip_size_t __size)78__attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) { 79 return (void *) __ockl_dm_alloc(__size); 80 } free(void * __ptr)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