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