[PATCH] D101106: [HIP] Fix overloaded function for _Float16

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added a reviewer: tra.
yaxunl requested review of this revision.

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


https://reviews.llvm.org/D101106

Files:
  clang/lib/Headers/__clang_hip_cmath.h
  clang/test/Headers/hip-header.hip


Index: clang/test/Headers/hip-header.hip
===================================================================
--- clang/test/Headers/hip-header.hip
+++ clang/test/Headers/hip-header.hip
@@ -25,12 +25,11 @@
   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);
 }
Index: clang/lib/Headers/__clang_hip_cmath.h
===================================================================
--- clang/lib/Headers/__clang_hip_cmath.h
+++ clang/lib/Headers/__clang_hip_cmath.h
@@ -258,7 +258,6 @@
   enum { value = 1 };
 };
 
-// ToDo: specializes is_arithmetic<_Float16>
 template <class _Tp> struct is_arithmetic {
   enum { value = 0 };
 };
@@ -307,6 +306,9 @@
 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;



Index: clang/test/Headers/hip-header.hip
===================================================================
--- clang/test/Headers/hip-header.hip
+++ clang/test/Headers/hip-header.hip
@@ -25,12 +25,11 @@
   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);
 }
Index: clang/lib/Headers/__clang_hip_cmath.h
===================================================================
--- clang/lib/Headers/__clang_hip_cmath.h
+++ clang/lib/Headers/__clang_hip_cmath.h
@@ -258,7 +258,6 @@
   enum { value = 1 };
 };
 
-// ToDo: specializes is_arithmetic<_Float16>
 template <class _Tp> struct is_arithmetic {
   enum { value = 0 };
 };
@@ -307,6 +306,9 @@
 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;

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits