166 lines · c
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-exception6 *7 *===-----------------------------------------------------------------------===8 */9 10/*11 * WARNING: This header is intended to be directly -include'd by12 * 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#define __cluster_dims__(...) __attribute__((cluster_dims(__VA_ARGS__)))29#define __no_cluster__ __attribute__((no_cluster))30 31#if !defined(__cplusplus) || __cplusplus < 201103L32 #define nullptr NULL;33#endif34 35#ifdef __cplusplus36extern "C" {37 __attribute__((__visibility__("default")))38 __attribute__((weak))39 __attribute__((noreturn))40 __device__ void __cxa_pure_virtual(void) {41 __builtin_trap();42 }43 __attribute__((__visibility__("default")))44 __attribute__((weak))45 __attribute__((noreturn))46 __device__ void __cxa_deleted_virtual(void) {47 __builtin_trap();48 }49}50#endif //__cplusplus51 52#if !defined(__HIPCC_RTC__)53#if __has_include("hip/hip_version.h")54#include "hip/hip_version.h"55#endif // __has_include("hip/hip_version.h")56#endif // __HIPCC_RTC__57 58typedef __SIZE_TYPE__ __hip_size_t;59 60#ifdef __cplusplus61extern "C" {62#endif //__cplusplus63 64#if HIP_VERSION_MAJOR * 100 + HIP_VERSION_MINOR >= 40565__device__ unsigned long long __ockl_dm_alloc(unsigned long long __size);66__device__ void __ockl_dm_dealloc(unsigned long long __addr);67#if __has_feature(address_sanitizer)68__device__ unsigned long long __asan_malloc_impl(unsigned long long __size,69 unsigned long long __pc);70__device__ void __asan_free_impl(unsigned long long __addr,71 unsigned long long __pc);72__attribute__((noinline, weak)) __device__ void *malloc(__hip_size_t __size) {73 unsigned long long __pc = (unsigned long long)__builtin_return_address(0);74 return (void *)__asan_malloc_impl(__size, __pc);75}76__attribute__((noinline, weak)) __device__ void free(void *__ptr) {77 unsigned long long __pc = (unsigned long long)__builtin_return_address(0);78 __asan_free_impl((unsigned long long)__ptr, __pc);79}80#else // __has_feature(address_sanitizer)81__attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {82 return (void *) __ockl_dm_alloc(__size);83}84__attribute__((weak)) inline __device__ void free(void *__ptr) {85 __ockl_dm_dealloc((unsigned long long)__ptr);86}87#endif // __has_feature(address_sanitizer)88#else // HIP version check89#if __HIP_ENABLE_DEVICE_MALLOC__90__device__ void *__hip_malloc(__hip_size_t __size);91__device__ void *__hip_free(void *__ptr);92__attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {93 return __hip_malloc(__size);94}95__attribute__((weak)) inline __device__ void free(void *__ptr) {96 __hip_free(__ptr);97}98#else // __HIP_ENABLE_DEVICE_MALLOC__99__attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {100 __builtin_trap();101 return (void *)0;102}103__attribute__((weak)) inline __device__ void free(void *__ptr) {104 __builtin_trap();105}106#endif // __HIP_ENABLE_DEVICE_MALLOC__107#endif // HIP version check108 109#ifdef __cplusplus110} // extern "C"111#endif //__cplusplus112 113#if !defined(__HIPCC_RTC__)114#include <cmath>115#include <cstdlib>116#include <stdlib.h>117#if __has_include("hip/hip_version.h")118#include "hip/hip_version.h"119#endif // __has_include("hip/hip_version.h")120#else121typedef __SIZE_TYPE__ size_t;122// Define macros which are needed to declare HIP device API's without standard123// C/C++ headers. This is for readability so that these API's can be written124// the same way as non-hipRTC use case. These macros need to be popped so that125// they do not pollute users' name space.126#pragma push_macro("NULL")127#pragma push_macro("uint32_t")128#pragma push_macro("uint64_t")129#pragma push_macro("CHAR_BIT")130#pragma push_macro("INT_MAX")131#pragma push_macro("INT_MIN")132#define NULL (void *)0133#define uint32_t __UINT32_TYPE__134#define uint64_t __UINT64_TYPE__135#define CHAR_BIT __CHAR_BIT__136#define INT_MAX __INTMAX_MAX__137#define INT_MIN (-__INT_MAX__ - 1)138#endif // __HIPCC_RTC__139 140#include <__clang_hip_libdevice_declares.h>141#include <__clang_hip_math.h>142#include <__clang_hip_stdlib.h>143 144#if defined(__HIPCC_RTC__)145#include <__clang_hip_cmath.h>146#else147#include <__clang_cuda_math_forward_declares.h>148#include <__clang_hip_cmath.h>149#include <__clang_cuda_complex_builtins.h>150#include <algorithm>151#include <complex>152#include <new>153#endif // __HIPCC_RTC__154 155#define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1156#if defined(__HIPCC_RTC__)157#pragma pop_macro("NULL")158#pragma pop_macro("uint32_t")159#pragma pop_macro("uint64_t")160#pragma pop_macro("CHAR_BIT")161#pragma pop_macro("INT_MAX")162#pragma pop_macro("INT_MIN")163#endif // __HIPCC_RTC__164#endif // __HIP__165#endif // __CLANG_HIP_RUNTIME_WRAPPER_H__166