Make hip math headers easier to use from C

[JonChesterfield]

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.

[jdoerfert]

I'm fine with this @yaxunl should accept though.

clang/lib/Headers/__clang_hip_math.h

561

Nit: You mix the C and C++ math declarations in this file, while possible, I somehow thing the cuda_{cmath/math} split is nicer.

[yaxunl]

clang/lib/Headers/__clang_hip_math.h

98

typo ?

558

typo

561

right

1185

typo

[JonChesterfield]

• Fix missing underscores

[JonChesterfield]

clang/lib/Headers/__clang_hip_math.h

561

Open to me implementing that in a later patch?

[yaxunl]

LGTM. Thanks.

clang/lib/Headers/__clang_hip_math.h

561

If you could work on it that would be great.

clang/lib/Headers/__clang_hip_libdevice_declares.h

 /*===---- __clang_hip_libdevice_declares.h - HIP device library decls -------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #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);
 __device__ __attribute__((pure)) float __ocml_acosh_f32(float);
 __device__ __attribute__((const)) float __ocml_asin_f32(float);
 __device__ __attribute__((pure)) float __ocml_asinh_f32(float);
 __device__ __attribute__((const)) float __ocml_atan2_f32(float, float);
 __device__ __attribute__((const)) float __ocml_atan_f32(float);

 __device__ __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
 __device__ __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
 __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);
 __device__ __2f16 __ocml_sin_2f16(__2f16);
 __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__

clang/lib/Headers/__clang_hip_math.h

if (*__tagp == '0') {
     else
       return __make_mantissa_base8(__tagp);
   }
 
   return __make_mantissa_base10(__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__
 inline float acoshf(float __x) { return __ocml_acosh_f32(__x); }
 __DEVICE__
 inline float asinf(float __x) { return __ocml_asin_f32(__x); }
 __DEVICE__
 inline float asinhf(float __x) { return __ocml_asinh_f32(__x); }
union {
     float val;
     struct ieee_float {
       uint32_t mantissa : 22;
       uint32_t quiet : 1;
       uint32_t exponent : 8;
       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;
   __tmp.bits.exponent = ~0u;
   __tmp.bits.quiet = 1u;
   __tmp.bits.mantissa = __make_mantissa(__tagp);
 
   return __tmp.val;

 __DEVICE__
 inline float __sinf(float __x) { return __ocml_native_sin_f32(__x); }
 __DEVICE__
 inline float __tanf(float __x) { return __ocml_tan_f32(__x); }
 // END INTRINSICS
 // 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__
 inline double acosh(double __x) { return __ocml_acosh_f64(__x); }
 __DEVICE__
 inline double asin(double __x) { return __ocml_asin_f64(__x); }
 __DEVICE__
 inline double asinh(double __x) { return __ocml_asinh_f64(__x); }
#if !_WIN32
   union {
     double val;
     struct ieee_double {
       uint64_t mantissa : 51;
       uint32_t quiet : 1;
       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;
   __tmp.bits.exponent = ~0u;
   __tmp.bits.quiet = 1u;
   __tmp.bits.mantissa = __make_mantissa(__tagp);
 
   return __tmp.val;

 __host__ inline static int min(int __arg1, int __arg2) {
   return std::min(__arg1, __arg2);
 }
 
 __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); }
 
 __DEVICE__
 inline double pow(double __base, int __iexp) { return powi(__base, __iexp); }
 
 __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")
 #pragma pop_macro("__DEF_FUNI")
 #pragma pop_macro("__DEF_FLOAT_FUN2I")
 #pragma pop_macro("__HIP_OVERLOAD1")
 #pragma pop_macro("__HIP_OVERLOAD2")
 #pragma pop_macro("__DEVICE__")
 #pragma pop_macro("__RETURN_TYPE")
 
 #endif // __CLANG_HIP_MATH_H__