replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Make hip math headers easier to use from C 

Summary:
Make hip math headers easier to use from C

Motivation is a step towards using the hip math headers to implement math.h
for openmp, which needs to work with C as well as C++. NFC for C++ code.

Reviewers: yaxunl, jdoerfert

Reviewed By: yaxunl

Subscribers: sstefan1, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D84476
This commit is contained in:
Jon Chesterfield 2020-07-24 20:50:25 +01:00
parent 7d076e19e3
commit 679158e662
2 changed files with 13 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
clang/lib/Headers/__clang_hip_libdevice_declares.h
View File

@ -10,7 +10,9 @@
#ifndef __CLANG_HIP_LIBDEVICE_DECLARES_H__
#define __CLANG_HIP_LIBDEVICE_DECLARES_H__
#ifdef __cplusplus
extern "C" {
#endif
// BEGIN FLOAT
__device__ __attribute__((const)) float __ocml_acos_f32(float);
@ -316,7 +318,7 @@ __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
__device__ inline __2f16
__llvm_amdgcn_rcp_2f16(__2f16 __x) // Not currently exposed by ROCDL.
{
return __2f16{__llvm_amdgcn_rcp_f16(__x.x), __llvm_amdgcn_rcp_f16(__x.y)};
return (__2f16){__llvm_amdgcn_rcp_f16(__x.x), __llvm_amdgcn_rcp_f16(__x.y)};
}
__device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
__device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
@ -325,6 +327,8 @@ __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
__device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
__device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16);
#ifdef __cplusplus
} // extern "C"
#endif
#endif // __CLANG_HIP_LIBDEVICE_DECLARES_H__

10
clang/lib/Headers/__clang_hip_math.h
View File

@ -95,8 +95,10 @@ inline uint64_t __make_mantissa(const char *__tagp) {
}
// BEGIN FLOAT
#ifdef __cplusplus
__DEVICE__
inline float abs(float __x) { return __ocml_fabs_f32(__x); }
#endif
__DEVICE__
inline float acosf(float __x) { return __ocml_acos_f32(__x); }
__DEVICE__
@ -251,7 +253,7 @@ inline float nanf(const char *__tagp) {
uint32_t sign : 1;
} bits;
static_assert(sizeof(float) == sizeof(ieee_float), "");
static_assert(sizeof(float) == sizeof(struct ieee_float), "");
} __tmp;
__tmp.bits.sign = 0u;
@ -553,8 +555,10 @@ inline float __tanf(float __x) { return __ocml_tan_f32(__x); }
// END FLOAT
// BEGIN DOUBLE
#ifdef __cplusplus
__DEVICE__
inline double abs(double __x) { return __ocml_fabs_f64(__x); }
#endif
__DEVICE__
inline double acos(double __x) { return __ocml_acos_f64(__x); }
__DEVICE__
@ -712,7 +716,7 @@ inline double nan(const char *__tagp) {
uint32_t exponent : 11;
uint32_t sign : 1;
} bits;
static_assert(sizeof(double) == sizeof(ieee_double), "");
static_assert(sizeof(double) == sizeof(struct ieee_double), "");
} __tmp;
__tmp.bits.sign = 0u;
@ -1178,6 +1182,7 @@ __host__ inline static int max(int __arg1, int __arg2) {
return std::max(__arg1, __arg2);
}
#ifdef __cplusplus
__DEVICE__
inline float pow(float __base, int __iexp) { return powif(__base, __iexp); }
@ -1188,6 +1193,7 @@ __DEVICE__
inline _Float16 pow(_Float16 __base, int __iexp) {
return __ocml_pown_f16(__base, __iexp);
}
#endif
#pragma pop_macro("__DEF_FUN1")
#pragma pop_macro("__DEF_FUN2")