JonChesterfield created this revision. JonChesterfield added a reviewer: yaxunl. Herald added a reviewer: jdoerfert. Herald added subscribers: cfe-commits, sstefan1. Herald added a project: clang.
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. Repository: rG LLVM Github Monorepo https://reviews.llvm.org/D84476 Files: clang/lib/Headers/__clang_hip_libdevice_declares.h clang/lib/Headers/__clang_hip_math.h Index: clang/lib/Headers/__clang_hip_math.h =================================================================== --- clang/lib/Headers/__clang_hip_math.h +++ clang/lib/Headers/__clang_hip_math.h @@ -95,8 +95,10 @@ } // 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 @@ 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 @@ // 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 @@ 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 @@ return std::max(__arg1, __arg2); } +#ifdef _cplusplus __DEVICE__ inline float pow(float __base, int __iexp) { return powif(__base, __iexp); } @@ -1188,6 +1193,7 @@ inline _Float16 pow(_Float16 __base, int __iexp) { return __ocml_pown_f16(__base, __iexp); } +#endif #pragma pop_macro("__DEF_FUN1") #pragma pop_macro("__DEF_FUN2") Index: clang/lib/Headers/__clang_hip_libdevice_declares.h =================================================================== --- clang/lib/Headers/__clang_hip_libdevice_declares.h +++ clang/lib/Headers/__clang_hip_libdevice_declares.h @@ -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__ 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_trunc_2f16(__2f16); __device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16); +#ifdef __cplusplus } // extern "C" +#endif #endif // __CLANG_HIP_LIBDEVICE_DECLARES_H__
Index: clang/lib/Headers/__clang_hip_math.h =================================================================== --- clang/lib/Headers/__clang_hip_math.h +++ clang/lib/Headers/__clang_hip_math.h @@ -95,8 +95,10 @@ } // 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 @@ 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 @@ // 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 @@ 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 @@ return std::max(__arg1, __arg2); } +#ifdef _cplusplus __DEVICE__ inline float pow(float __base, int __iexp) { return powif(__base, __iexp); } @@ -1188,6 +1193,7 @@ inline _Float16 pow(_Float16 __base, int __iexp) { return __ocml_pown_f16(__base, __iexp); } +#endif #pragma pop_macro("__DEF_FUN1") #pragma pop_macro("__DEF_FUN2") Index: clang/lib/Headers/__clang_hip_libdevice_declares.h =================================================================== --- clang/lib/Headers/__clang_hip_libdevice_declares.h +++ clang/lib/Headers/__clang_hip_libdevice_declares.h @@ -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__ 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_trunc_2f16(__2f16); __device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16); +#ifdef __cplusplus } // extern "C" +#endif #endif // __CLANG_HIP_LIBDEVICE_DECLARES_H__
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits