@@ -65,6 +65,12 @@ The purpose of conversion from float to bfloat16 is to reduce ammount of memory
65
65
required to store floating-point numbers. Computations are expected to be done with
66
66
32-bit floating-point values.
67
67
68
+ This extension is an optional kernel feature as described in
69
+ https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:optional-kernel-features[section 5.7]
70
+ of the SYCL 2020 spec. Therefore, attempting to submit a kernel using this
71
+ feature to a device that does not support it should cause a synchronous
72
+ `errc::kernel_not_supported` exception to be thrown from the kernel invocation`
73
+ command (e.g. from `parallel_for`).
68
74
69
75
== Feature test macro
70
76
@@ -97,12 +103,14 @@ enum class aspect {
97
103
If a SYCL device has the `ext_intel_bf16_conversion` aspect, then it natively
98
104
supports conversion of values of `float` type to `bfloat16` and back.
99
105
100
- If the device doesn't have the aspect, objects of `bfloat16` class ust not be
106
+ If the device doesn't have the aspect, objects of `bfloat16` class must not be
101
107
used in the device code.
102
108
103
109
== New `bfloat16` class
104
110
105
- The following class provides the conversion functionality:
111
+ The `bfloat16` class below provides the conversion functionality. Conversion
112
+ from `float` to `bfloat16` is done with round to nearest even(RTE) rounding
113
+ mode.
106
114
107
115
[source]
108
116
----
@@ -117,11 +125,15 @@ bfloat16 {
117
125
storage_t value;
118
126
119
127
public:
128
+ bfloat16() = default;
129
+ bfloat16(const bfloat16&) = default;
130
+ ~bfloat16() = default;
131
+
120
132
// Direct initialization
121
- bfloat16(const storage_t& a )
133
+ bfloat16(const storage_t&)
122
134
123
135
// Convert from float to bfloat16
124
- bfloat16(const float& a );
136
+ bfloat16(const float&);
125
137
126
138
// Convert from bfloat16 to float
127
139
operator float() const;
@@ -140,6 +152,11 @@ public:
140
152
141
153
[source]
142
154
----
155
+ #include <sycl/sycl.hpp>
156
+ #include <sycl/ext/intel/experimental/bfloat16.hpp>
157
+
158
+ using sycl::ext::intel::experimental::bfloat16;
159
+
143
160
bfloat16 operator+(const bfloat16 &lhs, const bfloat16 &rhs) {
144
161
return static_cast<float>(lhs) + static_cast<float>(rhs);
145
162
}
@@ -150,7 +167,7 @@ float foo(float a, float b) {
150
167
bfloat16 B {b};
151
168
152
169
// Convert A and B from bfloat16 to float, do addition on floating-pointer
153
- // numbers, then convert the result to bfloat16 and store in C.
170
+ // numbers, then convert the result to bfloat16 and store it in C.
154
171
bfloat16 C = A + B;
155
172
156
173
// Return the result converted from bfloat16 to float.
0 commit comments