|
| 1 | += SYCL_INTEL_bf16_conversion |
| 2 | + |
| 3 | +:source-highlighter: coderay |
| 4 | +:coderay-linenums-mode: table |
| 5 | + |
| 6 | +// This section needs to be after the document title. |
| 7 | +:doctype: book |
| 8 | +:toc2: |
| 9 | +:toc: left |
| 10 | +:encoding: utf-8 |
| 11 | +:lang: en |
| 12 | + |
| 13 | +:blank: pass:[ +] |
| 14 | + |
| 15 | +// Set the default source code type in this document to C++, |
| 16 | +// for syntax highlighting purposes. This is needed because |
| 17 | +// docbook uses c++ and html5 uses cpp. |
| 18 | +:language: {basebackend@docbook:c++:cpp} |
| 19 | + |
| 20 | +// This is necessary for asciidoc, but not for asciidoctor |
| 21 | +:cpp: C++ |
| 22 | + |
| 23 | +== Notice |
| 24 | + |
| 25 | +IMPORTANT: This specification is a draft. |
| 26 | + |
| 27 | +Copyright (c) 2021 Intel Corporation. All rights reserved. |
| 28 | + |
| 29 | +NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are |
| 30 | +trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. |
| 31 | +used by permission by Khronos. |
| 32 | + |
| 33 | +NOTE: This document is better viewed when rendered as html with asciidoctor. |
| 34 | +GitHub does not render image icons. |
| 35 | + |
| 36 | +== Dependencies |
| 37 | + |
| 38 | +This extension is written against the SYCL 2020 specification, Revision 3. |
| 39 | + |
| 40 | +== Status |
| 41 | + |
| 42 | +Draft |
| 43 | + |
| 44 | +This is a preview extension specification, intended to provide early access to |
| 45 | +a feature for review and community feedback. When the feature matures, this |
| 46 | +specification may be released as a formal extension. |
| 47 | + |
| 48 | +Because the interfaces defined by this specification are not final and are |
| 49 | +subject to change they are not intended to be used by shipping software |
| 50 | +products. |
| 51 | + |
| 52 | +== Version |
| 53 | + |
| 54 | +Built On: {docdate} + |
| 55 | +Revision: 1 |
| 56 | + |
| 57 | +== Introduction |
| 58 | + |
| 59 | +This extension adds functionality to convert value of single-precision |
| 60 | +floating-point type(`float`) to `bfloat16` type and vice versa. The extension |
| 61 | +doesn't add support for `bfloat16` type as such, instead it uses 16-bit integer |
| 62 | +type(`uint16_t`) as a storage for `bfloat16` values. |
| 63 | + |
| 64 | +The purpose of conversion from float to bfloat16 is to reduce ammount of memory |
| 65 | +required to store floating-point numbers. Computations are expected to be done with |
| 66 | +32-bit floating-point values. |
| 67 | + |
| 68 | + |
| 69 | +== Feature test macro |
| 70 | + |
| 71 | +This extension provides a feature-test macro as described in the core SYCL |
| 72 | +specification section 6.3.3 "Feature test macros". Therefore, an implementation |
| 73 | +supporting this extension must predefine the macro |
| 74 | +`SYCL_EXT_INTEL_BF16_CONVERSION` to one of the values defined in the table |
| 75 | +below. Applications can test for the existence of this macro to determine if |
| 76 | +the implementation supports this feature, or applications can test the macro’s |
| 77 | + value to determine which of the extension’s APIs the implementation supports. |
| 78 | + |
| 79 | +[%header,cols="1,5"] |
| 80 | +|=== |
| 81 | +|Value |Description |
| 82 | +|1 |Initial extension version. Base features are supported. |
| 83 | +|=== |
| 84 | + |
| 85 | +== Extension to `enum class aspect` |
| 86 | + |
| 87 | +[source] |
| 88 | +---- |
| 89 | +namespace sycl { |
| 90 | +enum class aspect { |
| 91 | + ... |
| 92 | + ext_intel_bf16_conversion |
| 93 | +} |
| 94 | +} |
| 95 | +---- |
| 96 | + |
| 97 | +If a SYCL device has the `ext_intel_bf16_conversion` aspect, then it natively |
| 98 | +supports conversion of values of `float` type to `bfloat16` and back. |
| 99 | + |
| 100 | +If the device doesn't have the aspect, objects of `bfloat16` class ust not be |
| 101 | +used in the device code. |
| 102 | + |
| 103 | +== New `bfloat16` class |
| 104 | + |
| 105 | +The following class provides the conversion functionality: |
| 106 | + |
| 107 | +[source] |
| 108 | +---- |
| 109 | +namespace sycl { |
| 110 | +namespace ext { |
| 111 | +namespace intel { |
| 112 | +namespace experimental { |
| 113 | +
|
| 114 | +class bfloat16 { |
| 115 | + using storage_t = uint16_t; |
| 116 | + storage_t value; |
| 117 | +
|
| 118 | +public: |
| 119 | + // Direct initialization |
| 120 | + bfloat16(const storage_t& a) |
| 121 | +
|
| 122 | + // Convert from float to bfloat16 |
| 123 | + bfloat16(const float& a); |
| 124 | +
|
| 125 | + // Convert from bfloat16 to float |
| 126 | + operator float() const; |
| 127 | +
|
| 128 | + // Get bfloat16 as uint16. |
| 129 | + operator storage_t() const; |
| 130 | +}; |
| 131 | +
|
| 132 | +} // namespace experimental |
| 133 | +} // namespace intel |
| 134 | +} // namespace ext |
| 135 | +} // namespace sycl |
| 136 | +---- |
| 137 | + |
| 138 | +== Example |
| 139 | + |
| 140 | +[source] |
| 141 | +---- |
| 142 | +bfloat16 operator+(const bfloat16 &lhs, const bfloat16 &rhs) { |
| 143 | + return static_cast<float>(lhs) + static_cast<float>(rhs); |
| 144 | +} |
| 145 | +
|
| 146 | +float foo(float a, float b) { |
| 147 | + // Convert from float to bfloat16. |
| 148 | + bfloat16 A {a}; |
| 149 | + bfloat16 B {b}; |
| 150 | +
|
| 151 | + // Convert A and B from bfloat16 to float, do addition on floating-pointer |
| 152 | + // numbers, then convert the result to bfloat16 and store in C. |
| 153 | + bfloat16 C = A + B; |
| 154 | +
|
| 155 | + // Return the result converted from bfloat16 to float. |
| 156 | + return C; |
| 157 | +} |
| 158 | +
|
| 159 | +int main (int argc, char *argv[]) |
| 160 | +{ |
| 161 | + float data[3] = {7.0, 8.1, 0.0}; |
| 162 | + cl::sycl::queue deviceQueue; |
| 163 | + cl::sycl::buffer<float, 1> buf {data, cl::sycl::range<1> {3}}; |
| 164 | +
|
| 165 | + deviceQueue.submit ([&] (cl::sycl::handler& cgh) |
| 166 | + { |
| 167 | + auto numbers = buf.get_access<cl::sycl::access::mode::read_write> (cgh); |
| 168 | + cgh.single_task<class simple_kernel> ([=] () |
| 169 | + { |
| 170 | + numbers[2] = foo(numbers[0], numbers[1]); |
| 171 | +
|
| 172 | + }); |
| 173 | + }); |
| 174 | + return 0; |
| 175 | +} |
| 176 | +---- |
| 177 | + |
| 178 | +== Issues |
| 179 | + |
| 180 | +None. |
| 181 | + |
| 182 | +== Revision History |
| 183 | + |
| 184 | +[cols="5,15,15,70"] |
| 185 | +[grid="rows"] |
| 186 | +[options="header"] |
| 187 | +|======================================== |
| 188 | +|Rev|Date|Author|Changes |
| 189 | +|1|2021-08-02|Alexey Sotkin |*Initial public working draft* |
| 190 | +|======================================== |
| 191 | + |
| 192 | +//************************************************************************ |
| 193 | +//Other formatting suggestions: |
| 194 | +// |
| 195 | +//* Use *bold* text for host APIs, or [source] syntax highlighting. |
| 196 | +//* Use +mono+ text for device APIs, or [source] syntax highlighting. |
| 197 | +//* Use +mono+ text for extension names, types, or enum values. |
| 198 | +//* Use _italics_ for parameters. |
| 199 | +//************************************************************************ |
| 200 | +
|
0 commit comments