Converting some operators to hidden friends

Also reorder member functions and improve the test.

Signed-off-by: Alexey Sotkin <[email protected]>

Author: AlexeySotkin

sycl/include/sycl/ext/intel/experimental/bfloat16.hpp

@@ -60,22 +60,45 @@ class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 {
   // Get raw bits representation of bfloat16
   operator storage_t() const { return value; }
 
-// Assignment operators overloading
+  // Logical operators (!,||,&&) are covered if we can cast to bool
+  explicit operator bool() { return to_float(value) != 0.0f; }
+
+  // Unary minus operator overloading
+  friend bfloat16 operator-(bfloat16 &lhs) {
+    return bfloat16{-to_float(lhs.value)};
+  }
+
+// Increment and decrement operators overloading
+#define OP(op)                                                                 \
+  friend bfloat16 &operator op(bfloat16 &lhs) {                                \
+    float f = to_float(lhs.value);                                             \
+    lhs.value = from_float(op f);                                              \
+    return lhs;                                                                \
+  }                                                                            \
+  friend bfloat16 operator op(bfloat16 &lhs, int) {                            \
+    bfloat16 old = lhs;                                                        \
+    operator op(lhs);                                                          \
+    return old;                                                                \
+  }
+  OP(++)
+  OP(--)
+#undef OP
+
+  // Assignment operators overloading
 #define OP(op)                                                                 \
-  friend bfloat16 operator op(bfloat16 &lhs, const bfloat16 &rhs) {            \
+  friend bfloat16 &operator op(bfloat16 &lhs, const bfloat16 &rhs) {           \
     float f = static_cast<float>(lhs);                                         \
     f op static_cast<float>(rhs);                                              \
     return lhs = f;                                                            \
   }                                                                            \
                                                                                \
   template <typename T>                                                        \
-  friend bfloat16 operator op(bfloat16 &lhs, const T &rhs) {                   \
+  friend bfloat16 &operator op(bfloat16 &lhs, const T &rhs) {                  \
     float f = static_cast<float>(lhs);                                         \
-                                                                               \
     f op static_cast<float>(rhs);                                              \
     return lhs = f;                                                            \
   }                                                                            \
-  template <typename T> friend T operator op(T &lhs, const bfloat16 &rhs) {    \
+  template <typename T> friend T &operator op(T &lhs, const bfloat16 &rhs) {   \
     float f = static_cast<float>(lhs);                                         \
     f op static_cast<float>(rhs);                                              \
     return lhs = f;                                                            \
@@ -86,28 +109,6 @@ class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 {
   OP(/=)
 #undef OP
 
-// Increment and decrement operators overloading
-#define OP(op)                                                                 \
-  bfloat16 &operator op() {                                                    \
-    float f = to_float(value);                                                 \
-    value = from_float(op f);                                                  \
-    return *this;                                                              \
-  }                                                                            \
-  bfloat16 operator op(int) {                                                  \
-    bfloat16 old = *this;                                                      \
-    operator op();                                                             \
-    return old;                                                                \
-  }
-  OP(++)
-  OP(--)
-#undef OP
-
-  // Unary minus operator overloading
-  bfloat16 operator-() { return bfloat16{-to_float(value)}; }
-
-  // Logical operators (!,||,&&) are covered if we can cast to bool
-  explicit operator bool() { return to_float(value) != 0.0f; }
-
 // Binary operators overloading
 #define OP(type, op)                                                           \
   friend type operator op(const bfloat16 &lhs, const bfloat16 &rhs) {          \
@@ -132,6 +133,7 @@ class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 {
   OP(bool, <=)
   OP(bool, >=)
 #undef OP
+
   // Bitwise(|,&,~,^), modulo(%) and shift(<<,>>) operations are not supported
   // for floating-point types.
 };

sycl/test/extensions/bfloat16.cpp

@@ -7,6 +7,7 @@ using sycl::ext::intel::experimental::bfloat16;
 
 SYCL_EXTERNAL uint16_t some_bf16_intrinsic(uint16_t x, uint16_t y);
 
+__attribute__((noinline))
 float op(float a, float b) {
   bfloat16 A {a};
 // CHECK: [[A:%.*]] = tail call spir_func zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float %a)
@@ -16,11 +17,11 @@ float op(float a, float b) {
 // CHECK: [[B:%.*]] = tail call spir_func zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float %b)
 // CHECK-NOT: fptoui
 
-  bfloat16 C = static_cast<float>(A) + static_cast<float>(B);
+  bfloat16 C = A + B;
 // CHECK: [[A_float:%.*]] = tail call spir_func float @_Z27__spirv_ConvertBF16ToFINTELt(i16 zeroext [[A]])
 // CHECK: [[B_float:%.*]] = tail call spir_func float @_Z27__spirv_ConvertBF16ToFINTELt(i16 zeroext [[B]])
-// CHECK: %add = fadd float [[A_float]], [[B_float]]
-// CHECK: [[C:%.*]] = tail call spir_func zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float %add)
+// CHECK: [[Add:%.*]] = fadd float [[A_float]], [[B_float]]
+// CHECK: [[C:%.*]] = tail call spir_func zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float [[Add]])
 // CHECK-NOT: uitofp
 // CHECK-NOT: fptoui