[HIP] Fix overloaded function for _Float16

[yaxunl]

Summary

Specializes __hip::is_arithmetic<_Float16> so that
overloaded functions with _Float16 arguments resolve
properly in a similar way as standard C++ header.

clang/lib/Headers/__clang_hip_cmath.h

 };
 template <> struct is_integral<long long> {
   enum { value = 1 };
 };
 template <> struct is_integral<unsigned long long> {
   enum { value = 1 };
 };
 
-// ToDo: specializes is_arithmetic<_Float16>
 template <class _Tp> struct is_arithmetic {
   enum { value = 0 };
 };
 template <> struct is_arithmetic<bool> {
   enum { value = 1 };
 };
 template <> struct is_arithmetic<char> {
   enum { value = 1 };
template <> struct is_arithmetic<unsigned long long> {
   enum { value = 1 };
 };
 template <> struct is_arithmetic<float> {
   enum { value = 1 };
 };
 template <> struct is_arithmetic<double> {
   enum { value = 1 };
 };
+template <> struct is_arithmetic<_Float16> {
+  enum { value = 1 };
+};
 
 struct true_type {
   static const __constant__ bool value = true;
 };
 struct false_type {
   static const __constant__ bool value = false;
 };
 

clang/test/Headers/hip-header.hip

 };
 
 #if __cplusplus >= 201103L
 // Check __hip::__numeric_type can be used with a class without default ctor.
 __device__ void test_numeric_type() {
   int x = __hip::__numeric_type<Number>::value;
 }
 
-// ToDo: Fix __clang_hip_cmake.h to specialize __hip::is_arithmetic<_Float16>
-// to resolve fma(_Float16, _Float16, int) to fma(double, double, double)
+// Check fma(_Float16, _Float16, int) is resolved to fma(double, double, double)
 // instead of fma(_Float16, _Float16, _Float16).
 
 // CXX14-LABEL: define{{.*}}@_Z8test_fma
-// CXX14: call {{.*}}@__ocml_fma_f16
+// CXX14: call {{.*}}@__ocml_fma_f64
 __device__ double test_fma(_Float16 h, int i) {
   return fma(h, h, i);
 }
 
 #endif
 
 // CHECK-LABEL: amdgpu_kernel void @_Z4kernPff
 __global__ void kern(float *x, float y) {