[SYCL] Float2Half precision requirements, convert method correction.

garimagu · bader · commit 6a2cd906fb58 · 2019-08-30T00:02:38.000+03:00
- Added the support for output operator for half datatypes.
- Edit test case vec_convert.cpp for half datatype.
  Corrected source code for float2Half conversion.
- Correct the automatic/rte mode implementation.

Signed-off-by: Garima Gupta &lt;garima.gupta@intel.com&gt;
diff --git a/sycl/include/CL/sycl/half_type.hpp b/sycl/include/CL/sycl/half_type.hpp
@@ -10,6 +10,7 @@
 
 #include <cstdint>
 #include <functional>
+#include <iostream>
 
 namespace cl {
 namespace sycl {
@@ -95,3 +96,21 @@ template <> struct hash<cl::sycl::detail::half_impl::half> {
 };
 
 } // namespace std
+
+#ifdef __SYCL_DEVICE_ONLY__
+using half = _Float16;
+#else
+using half = cl::sycl::detail::half_impl::half;
+#endif
+
+inline std::ostream &operator<<(std::ostream &O, half const &rhs) {
+  O << static_cast<float>(rhs);
+  return O;
+}
+
+inline std::istream &operator>>(std::istream &I, half &rhs) {
+  float ValFloat = 0.0f;
+  I >> ValFloat;
+  rhs = ValFloat;
+  return I;
+}
diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp
@@ -51,16 +51,14 @@
 
 #include <array>
 #include <cmath>
+#ifndef __SYCL_DEVICE_ONLY__
+#include <cfenv>
+#pragma STDC FENV_ACCESS ON
+#endif
 
 // 4.10.1: Scalar data types
 // 4.10.2: SYCL vector types
 
-#ifdef __SYCL_DEVICE_ONLY__
-using half = _Float16;
-#else
-using half = cl::sycl::detail::half_impl::half;
-#endif
-
 namespace cl {
 namespace sycl {
 
@@ -258,6 +256,8 @@ detail::enable_if_t<std::is_same<T, R>::value, R> convertImpl(T Value) {
   return Value;
 }
 
+// Note for float to half conversions, static_cast calls the conversion operator
+// implemented for host that takes care of the precision requirements.
 template <typename T, typename R, rounding_mode roundingMode>
 detail::enable_if_t<!std::is_same<T, R>::value &&
                         (is_int_to_int<T, R>::value ||
@@ -270,16 +270,23 @@ convertImpl(T Value) {
 
 // float to int
 template <typename T, typename R, rounding_mode roundingMode>
-detail::enable_if_t<!std::is_same<T, R>::value && is_float_to_int<T, R>::value,
-                    R>
-convertImpl(T Value) {
+detail::enable_if_t<is_float_to_int<T, R>::value, R> convertImpl(T Value) {
 #ifndef __SYCL_DEVICE_ONLY__
   switch (roundingMode) {
     // Round to nearest even is default rounding mode for floating-point types
   case rounding_mode::automatic:
     // Round to nearest even.
-  case rounding_mode::rte:
-    return std::round(Value);
+  case rounding_mode::rte: {
+    int OldRoundingDirection = std::fegetround();
+    int Err = std::fesetround(FE_TONEAREST);
+    if (Err)
+      throw runtime_error("Unable to set rounding mode to FE_TONEAREST");
+    R Result = std::rint(Value);
+    Err = std::fesetround(OldRoundingDirection);
+    if (Err)
+      throw runtime_error("Unable to restore rounding mode.");
+    return Result;
+  }
     // Round toward zero.
   case rounding_mode::rtz:
     return std::trunc(Value);
@@ -294,7 +301,7 @@ convertImpl(T Value) {
     return static_cast<R>(Value);
   };
 #else
-  // TODO implement device side convertion.
+  // TODO implement device side conversion.
   return static_cast<R>(Value);
 #endif
 }
diff --git a/sycl/source/half_type.cpp b/sycl/source/half_type.cpp
@@ -9,7 +9,6 @@
 #include <CL/sycl/half_type.hpp>
 // This is included to enable __builtin_expect()
 #include <CL/sycl/detail/platform_util.hpp>
-#include <iostream>
 #include <cstring>
 
 namespace cl {
@@ -28,7 +27,13 @@ static uint16_t float2Half(const float &Val) {
   const int8_t Exp32Diff = Exp32 - 127;
 
   uint16_t Exp16 = 0;
+
+  // convert 23-bit mantissa to 10-bit mantissa.
   uint16_t Frac16 = Frac32 >> 13;
+  // Round the mantissa as given in OpenCL spec section : 6.1.1.1 The half data
+  // type.
+  if (Frac32 >> 12 & 0x01)
+    Frac16 += 1;
 
   if (__builtin_expect(Exp32 == 0xff || Exp32Diff > 15, 0)) {
     Exp16 = 0x1f;
@@ -111,18 +116,6 @@ static float half2Float(const uint16_t &Val) {
   return Result;
 }
 
-std::ostream &operator<<(std::ostream &O, const half_impl::half &Val) {
-  O << static_cast<float>(Val);
-  return O;
-}
-
-std::istream &operator>>(std::istream &I, half_impl::half &ValHalf) {
-  float ValFloat = 0.0f;
-  I >> ValFloat;
-  ValHalf = ValFloat;
-  return I;
-}
-
 namespace half_impl {
 
 half::half(const float &RHS) : Buf(float2Half(RHS)) {}
diff --git a/sycl/test/basic_tests/vec_convert.cpp b/sycl/test/basic_tests/vec_convert.cpp
@@ -71,27 +71,33 @@ int main() {
       int8{2, 3, 3, -2, -3, -3, 0, 0});
   test<float, int, 8, rounding_mode::automatic>(
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
-      int8{2, 3, 3, -2, -3, -3, 0, 0});
+      int8{2, 2, 3, -2, -2, -3, 0, 0});
   test<int, float, 8, rounding_mode::automatic>(
       int8{2, 3, 3, -2, -3, -3, 0, 0},
       float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f});
   test<float, float, 8, rounding_mode::automatic>(
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
+  test<float, half, 8, rounding_mode::automatic>(
+      float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
+      half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
 
   // rte
   test<int, int, 8, rounding_mode::rte>(
       int8{2, 3, 3, -2, -3, -3, 0, 0},
       int8{2, 3, 3, -2, -3, -3, 0, 0});
   test<float, int, 8, rounding_mode::rte>(
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
-      int8{2, 3, 3, -2, -3, -3, 0, 0});
+      int8{2, 2, 3, -2, -2, -3, 0, 0});
   test<int, float, 8, rounding_mode::rte>(
       int8{2, 3, 3, -2, -3, -3, 0, 0},
       float8{2.f, 3.f, 3.f, -2.f, -3.f, -3.f, 0.f, 0.f});
   test<float, float, 8, rounding_mode::rte>(
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
+  test<float, half, 8, rounding_mode::rte>(
+      float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
+      half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
 
   // rtz
   test<int, int, 8, rounding_mode::rtz>(
@@ -106,6 +112,9 @@ int main() {
   test<float, float, 8, rounding_mode::rtz>(
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
+  test<float, half, 8, rounding_mode::rtz>(
+      float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
+      half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
 
   // rtp
   test<int, int, 8, rounding_mode::rtp>(
@@ -120,6 +129,9 @@ int main() {
   test<float, float, 8, rounding_mode::rtp>(
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
+  test<float, half, 8, rounding_mode::rtp>(
+      float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
+      half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
 
   // rtn
   test<int, int, 8, rounding_mode::rtn>(
@@ -134,6 +146,9 @@ int main() {
   test<float, float, 8, rounding_mode::rtn>(
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
       float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
+  test<float, half, 8, rounding_mode::rtn>(
+      float8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f},
+      half8{+2.3f, +2.5f, +2.7f, -2.3f, -2.5f, -2.7f, 0.f, 0.f});
 
   return 0;
 }
