blob: 73021d256cbae90e8dc57db2fb14edebb43495ef [file] [log] [blame]
Sasha Smundak0fc590b2020-10-07 08:11:59 -07001/*===---- __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
Pirama Arumuga Nainar494f6452021-12-02 10:42:14 -080021#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
33extern "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
Pirama Arumuga Nainar7e1f8392021-08-16 17:30:48 -070049#if !defined(__HIPCC_RTC__)
Sasha Smundak0fc590b2020-10-07 08:11:59 -070050#include <cmath>
51#include <cstdlib>
52#include <stdlib.h>
Pirama Arumuga Nainar7e1f8392021-08-16 17:30:48 -070053#else
54typedef __SIZE_TYPE__ size_t;
55// Define macros which are needed to declare HIP device API's without standard
56// C/C++ headers. This is for readability so that these API's can be written
57// the same way as non-hipRTC use case. These macros need to be popped so that
58// they do not pollute users' name space.
59#pragma push_macro("NULL")
60#pragma push_macro("uint32_t")
61#pragma push_macro("uint64_t")
62#pragma push_macro("CHAR_BIT")
63#pragma push_macro("INT_MAX")
64#define NULL (void *)0
65#define uint32_t __UINT32_TYPE__
66#define uint64_t __UINT64_TYPE__
67#define CHAR_BIT __CHAR_BIT__
68#define INT_MAX __INTMAX_MAX__
69#endif // __HIPCC_RTC__
Sasha Smundak0fc590b2020-10-07 08:11:59 -070070
Pirama Arumuga Nainar494f6452021-12-02 10:42:14 -080071typedef __SIZE_TYPE__ __hip_size_t;
Sasha Smundak0fc590b2020-10-07 08:11:59 -070072
Pirama Arumuga Nainar494f6452021-12-02 10:42:14 -080073#ifdef __cplusplus
74extern "C" {
75#endif //__cplusplus
Sasha Smundak4b1f33a2021-01-11 15:05:07 -080076
Sasha Smundak0fc590b2020-10-07 08:11:59 -070077#if __HIP_ENABLE_DEVICE_MALLOC__
Pirama Arumuga Nainar494f6452021-12-02 10:42:14 -080078__device__ void *__hip_malloc(__hip_size_t __size);
79__device__ void *__hip_free(void *__ptr);
80__attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {
Sasha Smundak0fc590b2020-10-07 08:11:59 -070081 return __hip_malloc(__size);
82}
Pirama Arumuga Nainar494f6452021-12-02 10:42:14 -080083__attribute__((weak)) inline __device__ void *free(void *__ptr) {
84 return __hip_free(__ptr);
Sasha Smundak0fc590b2020-10-07 08:11:59 -070085}
Pirama Arumuga Nainar494f6452021-12-02 10:42:14 -080086#else
87__attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) {
Sasha Smundak0fc590b2020-10-07 08:11:59 -070088 __builtin_trap();
Pirama Arumuga Nainar494f6452021-12-02 10:42:14 -080089 return (void *)0;
90}
91__attribute__((weak)) inline __device__ void *free(void *__ptr) {
92 __builtin_trap();
93 return (void *)0;
Sasha Smundak0fc590b2020-10-07 08:11:59 -070094}
95#endif
96
Pirama Arumuga Nainar494f6452021-12-02 10:42:14 -080097#ifdef __cplusplus
98} // extern "C"
99#endif //__cplusplus
100
Sasha Smundak0fc590b2020-10-07 08:11:59 -0700101#include <__clang_hip_libdevice_declares.h>
102#include <__clang_hip_math.h>
103
Pirama Arumuga Nainar7e1f8392021-08-16 17:30:48 -0700104#if defined(__HIPCC_RTC__)
105#include <__clang_hip_cmath.h>
106#else
Sasha Smundak0fc590b2020-10-07 08:11:59 -0700107#include <__clang_cuda_math_forward_declares.h>
Sasha Smundak4b1f33a2021-01-11 15:05:07 -0800108#include <__clang_hip_cmath.h>
Sasha Smundak0fc590b2020-10-07 08:11:59 -0700109#include <__clang_cuda_complex_builtins.h>
Sasha Smundak0fc590b2020-10-07 08:11:59 -0700110#include <algorithm>
111#include <complex>
112#include <new>
Pirama Arumuga Nainar7e1f8392021-08-16 17:30:48 -0700113#endif // __HIPCC_RTC__
Sasha Smundak0fc590b2020-10-07 08:11:59 -0700114
115#define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1
Pirama Arumuga Nainar7e1f8392021-08-16 17:30:48 -0700116#if defined(__HIPCC_RTC__)
117#pragma pop_macro("NULL")
118#pragma pop_macro("uint32_t")
119#pragma pop_macro("uint64_t")
120#pragma pop_macro("CHAR_BIT")
121#pragma pop_macro("INT_MAX")
122#endif // __HIPCC_RTC__
Sasha Smundak0fc590b2020-10-07 08:11:59 -0700123#endif // __HIP__
124#endif // __CLANG_HIP_RUNTIME_WRAPPER_H__