clang/lib/Headers/__clang_hip_runtime_wrapper.h

   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__