Update SYCL_INTEL_bf16_conversion

Add explicit conversion functions
Add operator overloadings
Apply code review suggestions

Author: AlexeySotkin

sycl/doc/extensions/Bf16Conversion/SYCL_INTEL_bf16_conversion.asciidoc

@@ -52,7 +52,7 @@ products.
 == Version
 
 Built On: {docdate} +
-Revision: 1
+Revision: 2
 
 == Introduction
 
@@ -69,7 +69,7 @@ This extension is an optional kernel feature as described in
 https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:optional-kernel-features[section 5.7]
 of the SYCL 2020 spec. Therefore, attempting to submit a kernel using this
 feature to a device that does not support it should cause a synchronous
-`errc::kernel_not_supported` exception to be thrown from the kernel invocation`
+`errc::kernel_not_supported` exception to be thrown from the kernel invocation
 command (e.g. from `parallel_for`).
 
 == Feature test macro
@@ -119,27 +119,58 @@ namespace ext {
 namespace intel {
 namespace experimental {
 
-class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]]
 bfloat16 {
   using storage_t = uint16_t;
   storage_t value;
 
 public:
   bfloat16() = default;
-  bfloat16(const bfloat16&) = default;
+  bfloat16(const bfloat16 &) = default;
   ~bfloat16() = default;
 
+  // Explicit conversion functions
+  static storage_t from_float(const float &a);
+  static float to_float(const storage_t &a);
+
   // Direct initialization
-  bfloat16(const storage_t&)
+  bfloat16(const storage_t &a);
 
   // Convert from float to bfloat16
-  bfloat16(const float&);
+  bfloat16(const float &a);
 
   // Convert from bfloat16 to float
   operator float() const;
 
   // Get bfloat16 as uint16.
   operator storage_t() const;
+
+  // Convert to bool type
+  explicit operator bool()
+
+  friend bfloat16 operator-() { /* ... */ }
+
+  // OP is: prefix ++, --
+  friend bfloat16 &operatorOP() { /* ... */ }
+
+  // OP is: postfix ++, --
+  friend bfloat16 operatorOP(int) { /* ... */ }
+
+  // OP is: +=, -=, *=, /=
+  friend bfloat16 &operatorOP(bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ }
+
+  // OP is +, -, *, /
+  friend bfloat16 operatorOP(const bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ }
+  template <typename T>
+  friend bfloat16 operatorOP(const bfloat16 &lhs, const T &rhs) { /* ... */ }
+  template <typename T>
+  friend bfloat16 operatorOP(const T &lhs, const bfloat16 &rhs) { /* ... */ }
+
+  // OP is ==,!=, <, >, <=, >=
+  friend bool operatorOP(const bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ }
+  template <typename T>
+  friend bool operatorOP(const bfloat16 &lhs, const T &rhs) { /* ... */ }
+  template <typename T>
+  friend bool operatorOP(const T &lhs, const bfloat16 &rhs) { /* ... */ }
 };
 
 } // namespace experimental
@@ -148,6 +179,102 @@ public:
 } // namespace sycl
 ----
 
+Table 1. Member functions of `bfloat16` class.
+|===
+| Member Function | Description
+
+|  `static storage_t from_float(const float &a);`
+|  Explicitly convert from `float` to `bfloat16`.
+
+|  `static float to_float(const storage_t &a);`
+|  Interpret `a` as `bfloat16` and explicitly convert it to `float`.
+
+| `bfloat16(const storage_t& a)`
+| Initialize with an integer value which is interpreted as a `bfloat16` type.
+
+| `bfloat16(const float& a);`
+| Construct `bfloat16` from `float`. Converts `float` to `bfloat16`.
+
+| `operator float() const;`
+|  Return `bfloat16` value converted to `float`.
+
+| `operator storage_t() const;`
+| Return `uint16_t` value, whose bits represent `bfloat16` value.
+
+| `explicit operator bool() { /* ... */ }`
+| Convert `bfloat16` to `bool` type. Return `false` if the value equals to
+  zero, return `true` otherwise.
+
+| `friend bfloat16 operator-(bfloat16 &bf) { /* ... */ }`
+| Construct new instance of `bfloat16` class with negated value of the `bf`.
+
+| `friend bfloat16 &operatorOP(bfloat16 &bf) { /* ... */ }`
+| Perform an in-place `OP` prefix arithmetic operation on the `bf`,
+  assigning the result to the `bf` and return the `bf`.
+
+  OP is: `++, --`
+
+| `friend bfloat16 operatorOP(bfloat16 &bf, int) { /* ... */ }`
+| Perform an in-place `OP` postfix arithmetic operation on `bf`, assigning
+  the result to the `bf` and return a copy of `bf` before the operation is
+  performed.
+
+  OP is: `++, --`
+
+| `friend bfloat16 operatorOP(const bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ }`
+| Perform an in-place `OP` arithmetic operation between the `lhs` and the `rhs`
+  and return the `lhs`.
+
+  OP is: `+=, -=, *=, /=`
+
+| `friend type operatorOP(const bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ }`
+| Construct a new instance of the `bfloat16` class with the value of the new
+  `bfloat16` instance being the result of an OP arithmetic operation between
+  the `lhs` `bfloat16` and `rhs` `bfloat16` values.
+
+  OP is `+, -, *, /`
+
+| `template <typename T>
+  friend bfloat16 operatorOP(const bfloat16 &lhs, const T &rhs) { /* ... */ }`
+| Construct a new instance of the `bfloat16` class with the value of the new
+  `bfloat16` instance being the result of an OP arithmetic operation between
+  the `lhs` `bfloat16` value and `rhs` of template type `T`. Type `T` must be
+  convertible to `float`.
+
+  OP is `+, -, *, /`
+
+| `template <typename T>
+  friend bfloat16 operatorOP(const T &lhs, const bfloat16 &rhs) { /* ... */ }`
+| Construct a new instance of the `bfloat16` class with the value of the new
+  `bfloat16` instance being the result of an OP arithmetic operation between
+  the `lhs` of template type `T` and `rhs` `bfloat16` value. Type `T` must be
+  convertible to `float`.
+
+  OP is `+, -, *, /`
+
+| `friend bool operatorOP(const bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ }`
+| Perform comparison operation OP between `lhs` `bfloat16` and `rhs` `bfloat16`
+  values and return the result as a boolean value.
+
+OP is `==, !=, <, >, <=, >=`
+
+| `template <typename T>
+  friend bool operatorOP(const bfloat16 &lhs, const T &rhs) { /* ... */ }`
+| Perform comparison operation OP between `lhs` `bfloat16` and `rhs` of
+  template type `T` and return the result as a boolean value. Type `T` must be
+  convertible to `float`.
+
+OP is `==, !=, <, >, <=, >=`
+
+| `template <typename T>
+  friend bool operatorOP(const T &lhs, const bfloat16 &rhs) { /* ... */ }`
+| Perform comparison operation OP between `lhs` of template type `T` and `rhs`
+  `bfloat16` value and return the result as a boolean value. Type `T` must be
+  convertible to `float`.
+
+OP is `==, !=, <, >, <=, >=`
+|===
+
 == Example
 
 [source]
@@ -176,13 +303,13 @@ float foo(float a, float b) {
 
 int main (int argc, char *argv[]) {
   float data[3] = {7.0, 8.1, 0.0};
-  sycl::device dev{sycl::default_selector{}};
+  sycl::device dev;
   sycl::queue deviceQueue{dev};
-  sycl::buffer<float, 1> buf {data, cl::sycl::range<1> {3}};
+  sycl::buffer<float, 1> buf {data, sycl::range<1> {3}};
 
-  if (dev.has(aspect::ext_intel_bf16_conversion)) {
-    deviceQueue.submit ([&] (cl::sycl::handler& cgh) {
-      auto numbers = buf.get_access<cl::sycl::access::mode::read_write> (cgh);
+  if (dev.has(sycl::aspect::ext_intel_bf16_conversion)) {
+    deviceQueue.submit ([&] (sycl::handler& cgh) {
+      auto numbers = buf.get_access<sycl::access::mode::read_write> (cgh);
       cgh.single_task<class simple_kernel> ([=] () {
         numbers[2] = foo(numbers[0], numbers[1]);
       });
@@ -203,15 +330,8 @@ None.
 [options="header"]
 |========================================
 |Rev|Date|Author|Changes
-|1|2021-08-02|Alexey Sotkin |*Initial public working draft*
+|1|2021-08-02|Alexey Sotkin |Initial public working draft
+|2|2021-08-17|Alexey Sotkin |Add explicit conversion functions +
+                             Add operator overloadings +
+                             Apply code review suggestions
 |========================================
-
-//************************************************************************
-//Other formatting suggestions:
-//
-//* Use *bold* text for host APIs, or [source] syntax highlighting.
-//* Use +mono+ text for device APIs, or [source] syntax highlighting.
-//* Use +mono+ text for extension names, types, or enum values.
-//* Use _italics_ for parameters.
-//************************************************************************
-