From 025cf7ecb4373e248841227d3724ad4f94062b8e Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Tue, 25 Jan 2022 14:06:25 +0000 Subject: [PATCH 01/11] Added bfloat16 support for cuda backend. Added bfloat16 in oneapi experimental namespace. Signed-off-by: jack.kirk --- .../sycl/ext/oneapi/experimental/bfloat16.hpp | 161 ++++++++++++++++++ 1 file changed, 161 insertions(+) create mode 100644 sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp diff --git a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp new file mode 100644 index 0000000000000..329094634d9ad --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp @@ -0,0 +1,161 @@ +//==--------- bfloat16.hpp ------- SYCL bfloat16 conversion ----------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace ext { +namespace oneapi { +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() = default; + + // Explicit conversion functions + static storage_t from_float(const float &a) { +#if defined(__SYCL_DEVICE_ONLY__) +#if defined(__NVPTX__) + return __nvvm_f2bf16_rn(a); +#else + return __spirv_ConvertFToBF16INTEL(a); +#endif +#else + throw exception{errc::feature_not_supported, + "Bfloat16 conversion is not supported on host device"}; +#endif + } + static float to_float(const storage_t &a) { +#if defined(__SYCL_DEVICE_ONLY__) +#if defined(__NVPTX__) + unsigned int y = a; + y = y << 16; + float *res = reinterpret_cast(&y); + return *res; +#else + return __spirv_ConvertBF16ToFINTEL(a); +#endif +#else + throw exception{errc::feature_not_supported, + "Bfloat16 conversion is not supported on host device"}; +#endif + } + + static bfloat16 from_bits(const storage_t &a) { + bfloat16 res; + res.value = a; + return res; + } + + // Implicit conversion from float to bfloat16 + bfloat16(const float &a) { value = from_float(a); } + + bfloat16 &operator=(const float &rhs) { + value = from_float(rhs); + return *this; + } + + // Implicit conversion from bfloat16 to float + operator float() const { return to_float(value); } + operator sycl::half() const { return to_float(value); } + + // Get raw bits representation of bfloat16 + storage_t raw() const { return value; } + + // 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) { \ + float f = static_cast(lhs); \ + f op static_cast(rhs); \ + return lhs = f; \ + } \ + template \ + friend bfloat16 &operator op(bfloat16 &lhs, const T &rhs) { \ + float f = static_cast(lhs); \ + f op static_cast(rhs); \ + return lhs = f; \ + } \ + template friend T &operator op(T &lhs, const bfloat16 &rhs) { \ + float f = static_cast(lhs); \ + f op static_cast(rhs); \ + return lhs = f; \ + } + OP(+=) + OP(-=) + OP(*=) + OP(/=) +#undef OP + +// Binary operators overloading +#define OP(type, op) \ + friend type operator op(const bfloat16 &lhs, const bfloat16 &rhs) { \ + return type{static_cast(lhs) op static_cast(rhs)}; \ + } \ + template \ + friend type operator op(const bfloat16 &lhs, const T &rhs) { \ + return type{static_cast(lhs) op static_cast(rhs)}; \ + } \ + template \ + friend type operator op(const T &lhs, const bfloat16 &rhs) { \ + return type{static_cast(lhs) op static_cast(rhs)}; \ + } + OP(bfloat16, +) + OP(bfloat16, -) + OP(bfloat16, *) + OP(bfloat16, /) + OP(bool, ==) + OP(bool, !=) + OP(bool, <) + OP(bool, >) + OP(bool, <=) + OP(bool, >=) +#undef OP + + // Bitwise(|,&,~,^), modulo(%) and shift(<<,>>) operations are not supported + // for floating-point types. +}; + +} // namespace experimental +} // namespace intel +} // namespace ext + +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) From 66b4e3344bc7a9e514d857d4931ba26ed192b3f9 Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Tue, 25 Jan 2022 14:13:58 +0000 Subject: [PATCH 02/11] deleted intel namespace bfloat16. --- .../sycl/ext/intel/experimental/bfloat16.hpp | 150 ------------------ 1 file changed, 150 deletions(-) delete mode 100644 sycl/include/sycl/ext/intel/experimental/bfloat16.hpp diff --git a/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp b/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp deleted file mode 100644 index 5a51f3746e225..0000000000000 --- a/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp +++ /dev/null @@ -1,150 +0,0 @@ -//==--------- bfloat16.hpp ------- SYCL bfloat16 conversion ----------------==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#pragma once - -#include -#include - -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -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() = default; - - // Explicit conversion functions - static storage_t from_float(const float &a) { -#if defined(__SYCL_DEVICE_ONLY__) - return __spirv_ConvertFToBF16INTEL(a); -#else - throw exception{errc::feature_not_supported, - "Bfloat16 conversion is not supported on host device"}; -#endif - } - static float to_float(const storage_t &a) { -#if defined(__SYCL_DEVICE_ONLY__) - return __spirv_ConvertBF16ToFINTEL(a); -#else - throw exception{errc::feature_not_supported, - "Bfloat16 conversion is not supported on host device"}; -#endif - } - - static bfloat16 from_bits(const storage_t &a) { - bfloat16 res; - res.value = a; - return res; - } - - // Implicit conversion from float to bfloat16 - bfloat16(const float &a) { value = from_float(a); } - - bfloat16 &operator=(const float &rhs) { - value = from_float(rhs); - return *this; - } - - // Implicit conversion from bfloat16 to float - operator float() const { return to_float(value); } - operator sycl::half() const { return to_float(value); } - - // Get raw bits representation of bfloat16 - storage_t raw() const { return value; } - - // 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) { \ - float f = static_cast(lhs); \ - f op static_cast(rhs); \ - return lhs = f; \ - } \ - template \ - friend bfloat16 &operator op(bfloat16 &lhs, const T &rhs) { \ - float f = static_cast(lhs); \ - f op static_cast(rhs); \ - return lhs = f; \ - } \ - template friend T &operator op(T &lhs, const bfloat16 &rhs) { \ - float f = static_cast(lhs); \ - f op static_cast(rhs); \ - return lhs = f; \ - } - OP(+=) - OP(-=) - OP(*=) - OP(/=) -#undef OP - -// Binary operators overloading -#define OP(type, op) \ - friend type operator op(const bfloat16 &lhs, const bfloat16 &rhs) { \ - return type{static_cast(lhs) op static_cast(rhs)}; \ - } \ - template \ - friend type operator op(const bfloat16 &lhs, const T &rhs) { \ - return type{static_cast(lhs) op static_cast(rhs)}; \ - } \ - template \ - friend type operator op(const T &lhs, const bfloat16 &rhs) { \ - return type{static_cast(lhs) op static_cast(rhs)}; \ - } - OP(bfloat16, +) - OP(bfloat16, -) - OP(bfloat16, *) - OP(bfloat16, /) - OP(bool, ==) - OP(bool, !=) - OP(bool, <) - OP(bool, >) - OP(bool, <=) - OP(bool, >=) -#undef OP - - // Bitwise(|,&,~,^), modulo(%) and shift(<<,>>) operations are not supported - // for floating-point types. -}; - -} // namespace experimental -} // namespace intel -} // namespace ext - -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) From 2d04406d0198b5321cf2aaa870d395e9f042755b Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Tue, 25 Jan 2022 14:29:32 +0000 Subject: [PATCH 03/11] Format. --- sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp index 329094634d9ad..ef1f01d5340ae 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp @@ -154,7 +154,7 @@ class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 { }; } // namespace experimental -} // namespace intel +} // namespace oneapi } // namespace ext } // namespace sycl From 9418f74ee1e1a35918f5bcf99a3d0d57f29dac90 Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Tue, 25 Jan 2022 14:35:02 +0000 Subject: [PATCH 04/11] Changed extension macro name. --- sycl/include/CL/sycl/feature_test.hpp.in | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/feature_test.hpp.in b/sycl/include/CL/sycl/feature_test.hpp.in index e6053ebf4ff1c..9bd849ca27d8a 100644 --- a/sycl/include/CL/sycl/feature_test.hpp.in +++ b/sycl/include/CL/sycl/feature_test.hpp.in @@ -46,7 +46,7 @@ namespace sycl { #define SYCL_EXT_ONEAPI_USE_PINNED_HOST_MEMORY_PROPERTY 1 #define SYCL_EXT_ONEAPI_SRGB 1 #define SYCL_EXT_ONEAPI_SUB_GROUP 1 -#define SYCL_EXT_INTEL_BF16_CONVERSION 1 +#define SYCL_EXT_ONEAPI_BF16_CONVERSION 1 #define SYCL_EXT_INTEL_BITCAST 1 #define SYCL_EXT_INTEL_DATAFLOW_PIPES 1 #ifdef __clang__ From 4d99f3f97c06dc529d1d6b16df645b1af27fa8a6 Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Thu, 17 Feb 2022 10:22:48 +0000 Subject: [PATCH 05/11] fixed test. --- sycl/test/extensions/bfloat16.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test/extensions/bfloat16.cpp b/sycl/test/extensions/bfloat16.cpp index 6be5459642d0c..dd87806942f8a 100644 --- a/sycl/test/extensions/bfloat16.cpp +++ b/sycl/test/extensions/bfloat16.cpp @@ -2,10 +2,10 @@ // UNSUPPORTED: cuda || hip_amd -#include +#include #include -using sycl::ext::intel::experimental::bfloat16; +using sycl::ext::oneapi::experimental::bfloat16; SYCL_EXTERNAL uint16_t some_bf16_intrinsic(uint16_t x, uint16_t y); SYCL_EXTERNAL void foo(long x, sycl::half y); From 3982001259745c617ad78d57dc67512e8d7ff6e9 Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Fri, 4 Mar 2022 15:41:39 +0000 Subject: [PATCH 06/11] Used neg ptx7.0 builtin for unary minus --- clang/include/clang/Basic/BuiltinsNVPTX.def | 5 +++++ llvm/include/llvm/IR/IntrinsicsNVVM.td | 13 +++++++++++++ llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 9 +++++++++ .../sycl/ext/oneapi/experimental/bfloat16.hpp | 13 +++++++++++-- 4 files changed, 38 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.def b/clang/include/clang/Basic/BuiltinsNVPTX.def index 449e4d1256944..955dbbaae8f0d 100644 --- a/clang/include/clang/Basic/BuiltinsNVPTX.def +++ b/clang/include/clang/Basic/BuiltinsNVPTX.def @@ -182,6 +182,11 @@ BUILTIN(__nvvm_fabs_ftz_f, "ff", "") BUILTIN(__nvvm_fabs_f, "ff", "") BUILTIN(__nvvm_fabs_d, "dd", "") +// Neg + +TARGET_BUILTIN(__nvvm_neg_bf16, "ZUsZUs", "", AND(SM_80,PTX70)) +TARGET_BUILTIN(__nvvm_neg_bf16x2, "ZUiZUi", "", AND(SM_80,PTX70)) + // Round BUILTIN(__nvvm_round_ftz_f, "ff", "") diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td index 33ba30d782ff3..b7b0813f05292 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -740,6 +740,19 @@ let TargetPrefix = "nvvm" in { def int_nvvm_fabs_d : GCCBuiltin<"__nvvm_fabs_d">, DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>; +// +// Neg bf16, bf16x2 +// + + foreach unary = ["neg"] in { + def int_nvvm_ # unary # _bf16 : + GCCBuiltin, + DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty], [IntrNoMem]>; + def int_nvvm_ # unary # _bf16x2 : + GCCBuiltin, + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]>; + } + // // Round // diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index ec004c5923ece..af9e1270bc5f5 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -719,6 +719,15 @@ def INT_NVVM_FABS_F : F_MATH_1<"abs.f32 \t$dst, $src0;", Float32Regs, def INT_NVVM_FABS_D : F_MATH_1<"abs.f64 \t$dst, $src0;", Float64Regs, Float64Regs, int_nvvm_fabs_d>; +// +// Neg bf16, bf16x2 +// + +def INT_NVVM_NEG_BF16 : F_MATH_1<"neg.bf16 \t$dst, $src0;", Int16Regs, + Int16Regs, int_nvvm_neg_bf16>; +def INT_NVVM_NEG_BF16X2 : F_MATH_1<"neg.bf16x2 \t$dst, $src0;", Int32Regs, + Int32Regs, int_nvvm_neg_bf16x2>; + // // Round // diff --git a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp index ef1f01d5340ae..3768c65aab6a3 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp @@ -42,7 +42,7 @@ class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 { static float to_float(const storage_t &a) { #if defined(__SYCL_DEVICE_ONLY__) #if defined(__NVPTX__) - unsigned int y = a; + uint32_t y = a; y = y << 16; float *res = reinterpret_cast(&y); return *res; @@ -81,7 +81,16 @@ class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 { // Unary minus operator overloading friend bfloat16 operator-(bfloat16 &lhs) { - return bfloat16{-to_float(lhs.value)}; +#if defined(__SYCL_DEVICE_ONLY__) +#if defined(__NVPTX__) + return from_bits(__nvvm_neg_bf16(lhs.value)); +#else + return bfloat16{-__spirv_ConvertBF16ToFINTEL(lhs.value)}; +#endif +#else + throw exception{errc::feature_not_supported, + "Bfloat16 unary minus is not supported on host device"}; +#endif } // Increment and decrement operators overloading From 8d2d11fecfbc1e9caebdc460c7a9cb19b4e18774 Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Mon, 7 Mar 2022 16:47:32 +0000 Subject: [PATCH 07/11] Replaced SYCL_EXT_INTEL_BF16_CONVERSION.asciidoc with SYCL_EXT_ONEAPI_BF16_CONVERSION.asciidoc --- .../SYCL_EXT_ONEAPI_BF16_CONVERSION.asciidoc | 336 ++++++++++++++++++ 1 file changed, 336 insertions(+) create mode 100644 sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_BF16_CONVERSION.asciidoc diff --git a/sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_BF16_CONVERSION.asciidoc b/sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_BF16_CONVERSION.asciidoc new file mode 100644 index 0000000000000..bf0a799671ffa --- /dev/null +++ b/sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_BF16_CONVERSION.asciidoc @@ -0,0 +1,336 @@ += sycl_oneapi_bf16_conversion + +:source-highlighter: coderay +:coderay-linenums-mode: table + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en + +:blank: pass:[ +] + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + +// This is necessary for asciidoc, but not for asciidoctor +:cpp: C++ + +== Notice + +IMPORTANT: This specification is a draft. + +Copyright (c) 2021-2022 Intel Corporation. All rights reserved. + +NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are +trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. +used by permission by Khronos. + +== Dependencies + +This extension is written against the SYCL 2020 specification, Revision 4. + +== Status + +Draft + +This is a preview extension specification, intended to provide early access to +a feature for review and community feedback. When the feature matures, this +specification may be released as a formal extension. + +Because the interfaces defined by this specification are not final and are +subject to change they are not intended to be used by shipping software +products. + +== Version + +Revision: 4 + +== Introduction + +This extension adds functionality to convert value of single-precision +floating-point type(`float`) to `bfloat16` type and vice versa. The extension +doesn't add support for `bfloat16` type as such, instead it uses 16-bit integer +type(`uint16_t`) as a storage for `bfloat16` values. + +The purpose of conversion from float to bfloat16 is to reduce ammount of memory +required to store floating-point numbers. Computations are expected to be done with +32-bit floating-point values. + +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 +command (e.g. from `parallel_for`). + +== Feature test macro + +This extension provides a feature-test macro as described in the core SYCL +specification section 6.3.3 "Feature test macros". Therefore, an implementation +supporting this extension must predefine the macro +`SYCL_EXT_ONEAPI_BF16_CONVERSION` to one of the values defined in the table +below. Applications can test for the existence of this macro to determine if +the implementation supports this feature, or applications can test the macro’s + value to determine which of the extension’s APIs the implementation supports. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version. Base features are supported. +|=== + +== Extension to `enum class aspect` + +[source] +---- +namespace sycl { +enum class aspect { + ... + ext_oneapi_bf16_conversion +} +} +---- + +If a SYCL device has the `ext_oneapi_bf16_conversion` aspect, then it natively +supports conversion of values of `float` type to `bfloat16` and back. + +If the device doesn't have the aspect, objects of `bfloat16` class must not be +used in the device code. + +**NOTE**: The `ext_oneapi_bf16_conversion` aspect is not yet supported. The +`bfloat16` class is currently supported only on Xe HP GPU and Nvidia A100 GPU. + +== New `bfloat16` class + +The `bfloat16` class below provides the conversion functionality. Conversion +from `float` to `bfloat16` is done with round to nearest even(RTE) rounding +mode. + +[source] +---- +namespace sycl { +namespace ext { +namespace oneapi { +namespace experimental { + +class bfloat16 { + using storage_t = uint16_t; + storage_t value; + +public: + 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); + + // Convert from float to bfloat16 + bfloat16(const float &a); + bfloat16 &operator=(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-(bfloat16 &bf) { /* ... */ } + + // OP is: prefix ++, -- + friend bfloat16 &operatorOP(bfloat16 &bf) { /* ... */ } + + // OP is: postfix ++, -- + friend bfloat16 operatorOP(bfloat16 &bf, int) { /* ... */ } + + // OP is: +=, -=, *=, /= + friend bfloat16 &operatorOP(bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ } + + // OP is +, -, *, / + friend bfloat16 operatorOP(const bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ } + template + friend bfloat16 operatorOP(const bfloat16 &lhs, const T &rhs) { /* ... */ } + template + friend bfloat16 operatorOP(const T &lhs, const bfloat16 &rhs) { /* ... */ } + + // OP is ==,!=, <, >, <=, >= + friend bool operatorOP(const bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ } + template + friend bool operatorOP(const bfloat16 &lhs, const T &rhs) { /* ... */ } + template + friend bool operatorOP(const T &lhs, const bfloat16 &rhs) { /* ... */ } +}; + +} // namespace experimental +} // namespace oneapi +} // namespace ext +} // 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 float& a);` +| Construct `bfloat16` from `float`. Converts `float` to `bfloat16`. + +| `bfloat16 &operator=(const float &a);` +| Replace the value with `a` converted 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 + 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 + 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 + 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 + 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] +---- +#include +#include + +using sycl::ext::oneapi::experimental::bfloat16; + +bfloat16 operator+(const bfloat16 &lhs, const bfloat16 &rhs) { + return static_cast(lhs) + static_cast(rhs); +} + +float foo(float a, float b) { + // Convert from float to bfloat16. + bfloat16 A {a}; + bfloat16 B {b}; + + // Convert A and B from bfloat16 to float, do addition on floating-pointer + // numbers, then convert the result to bfloat16 and store it in C. + bfloat16 C = A + B; + + // Return the result converted from bfloat16 to float. + return C; +} + +int main (int argc, char *argv[]) { + float data[3] = {7.0, 8.1, 0.0}; + sycl::device dev; + sycl::queue deviceQueue{dev}; + sycl::buffer buf {data, sycl::range<1> {3}}; + + if (dev.has(sycl::aspect::ext_oneapi_bf16_conversion)) { + deviceQueue.submit ([&] (sycl::handler& cgh) { + auto numbers = buf.get_access (cgh); + cgh.single_task ([=] () { + numbers[2] = foo(numbers[0], numbers[1]); + }); + }); + } + return 0; +} +---- + +== Issues + +None. + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|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 +|3|2021-08-18|Alexey Sotkin |Remove `uint16_t` constructor +|4|2022-03-07|Jack Kirk |Switch from Intel vendor specific to oneapi +|======================================== From 8a29c4412c06b1246bdcd0fa5954b70957211e36 Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Tue, 15 Mar 2022 14:41:54 +0000 Subject: [PATCH 08/11] Renamed extension to cover all bfloat16 funct. Removed aspect reference: can be added once the ext_oneapi_bfloat16 aspect is merged. --- ....asciidoc => sycl_ext_oneapi_bfloat16.asciidoc} | 14 +++++++------- .../sycl/ext/oneapi/experimental/bfloat16.hpp | 2 +- 2 files changed, 8 insertions(+), 8 deletions(-) rename sycl/doc/extensions/experimental/{SYCL_EXT_ONEAPI_BF16_CONVERSION.asciidoc => sycl_ext_oneapi_bfloat16.asciidoc} (96%) diff --git a/sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_BF16_CONVERSION.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16.asciidoc similarity index 96% rename from sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_BF16_CONVERSION.asciidoc rename to sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16.asciidoc index bf0a799671ffa..175219e23c47f 100644 --- a/sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_BF16_CONVERSION.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16.asciidoc @@ -1,4 +1,4 @@ -= sycl_oneapi_bf16_conversion += sycl_ext_oneapi_bfloat16 :source-highlighter: coderay :coderay-linenums-mode: table @@ -73,7 +73,7 @@ command (e.g. from `parallel_for`). This extension provides a feature-test macro as described in the core SYCL specification section 6.3.3 "Feature test macros". Therefore, an implementation supporting this extension must predefine the macro -`SYCL_EXT_ONEAPI_BF16_CONVERSION` to one of the values defined in the table +`SYCL_EXT_ONEAPI_BFLOAT16` to one of the values defined in the table below. Applications can test for the existence of this macro to determine if the implementation supports this feature, or applications can test the macro’s value to determine which of the extension’s APIs the implementation supports. @@ -91,18 +91,18 @@ the implementation supports this feature, or applications can test the macro’s namespace sycl { enum class aspect { ... - ext_oneapi_bf16_conversion + ext_oneapi_bfloat16 } } ---- -If a SYCL device has the `ext_oneapi_bf16_conversion` aspect, then it natively +If a SYCL device has the `ext_oneapi_bfloat16` aspect, then it natively supports conversion of values of `float` type to `bfloat16` and back. If the device doesn't have the aspect, objects of `bfloat16` class must not be used in the device code. -**NOTE**: The `ext_oneapi_bf16_conversion` aspect is not yet supported. The +**NOTE**: The `ext_oneapi_bfloat16` aspect is not yet supported. The `bfloat16` class is currently supported only on Xe HP GPU and Nvidia A100 GPU. == New `bfloat16` class @@ -304,7 +304,7 @@ int main (int argc, char *argv[]) { sycl::queue deviceQueue{dev}; sycl::buffer buf {data, sycl::range<1> {3}}; - if (dev.has(sycl::aspect::ext_oneapi_bf16_conversion)) { + if (dev.has(sycl::aspect::ext_oneapi_bfloat16)) { deviceQueue.submit ([&] (sycl::handler& cgh) { auto numbers = buf.get_access (cgh); cgh.single_task ([=] () { @@ -332,5 +332,5 @@ None. Add operator overloadings + Apply code review suggestions |3|2021-08-18|Alexey Sotkin |Remove `uint16_t` constructor -|4|2022-03-07|Jack Kirk |Switch from Intel vendor specific to oneapi +|4|2022-03-07|Aidan Belton and Jack Kirk |Switch from Intel vendor specific to oneapi |======================================== diff --git a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp index 3768c65aab6a3..1190c80631928 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp @@ -17,7 +17,7 @@ namespace ext { namespace oneapi { namespace experimental { -class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 { +class bfloat16 { using storage_t = uint16_t; storage_t value; From f1fba08b89c43795e68ca5e991520904fa556a8a Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Wed, 16 Mar 2022 11:08:25 +0000 Subject: [PATCH 09/11] Updated macro name --- sycl/include/CL/sycl/feature_test.hpp.in | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/feature_test.hpp.in b/sycl/include/CL/sycl/feature_test.hpp.in index 516464e1b9fe3..11a7afdaed8ed 100644 --- a/sycl/include/CL/sycl/feature_test.hpp.in +++ b/sycl/include/CL/sycl/feature_test.hpp.in @@ -55,7 +55,7 @@ namespace sycl { #define SYCL_EXT_ONEAPI_SUB_GROUP 1 #define SYCL_EXT_ONEAPI_PROPERTIES 1 #define SYCL_EXT_ONEAPI_NATIVE_MATH 1 -#define SYCL_EXT_ONEAPI_BF16_CONVERSION 1 +#define SYCL_EXT_ONEAPI_BFLOAT16 1 #define SYCL_EXT_INTEL_DATAFLOW_PIPES 1 #ifdef __clang__ #if __has_extension(sycl_extended_atomics) From 461ddb8206653bacae049e7c46b245bc1b2ac954 Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Wed, 16 Mar 2022 15:57:51 +0000 Subject: [PATCH 10/11] Removed old extension doc --- .../sycl_ext_intel_bf16_conversion.asciidoc | 335 ------------------ 1 file changed, 335 deletions(-) delete mode 100644 sycl/doc/extensions/experimental/sycl_ext_intel_bf16_conversion.asciidoc diff --git a/sycl/doc/extensions/experimental/sycl_ext_intel_bf16_conversion.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_intel_bf16_conversion.asciidoc deleted file mode 100644 index 9b1018ced0b34..0000000000000 --- a/sycl/doc/extensions/experimental/sycl_ext_intel_bf16_conversion.asciidoc +++ /dev/null @@ -1,335 +0,0 @@ -= SYCL_INTEL_bf16_conversion - -:source-highlighter: coderay -:coderay-linenums-mode: table - -// This section needs to be after the document title. -:doctype: book -:toc2: -:toc: left -:encoding: utf-8 -:lang: en - -:blank: pass:[ +] - -// Set the default source code type in this document to C++, -// for syntax highlighting purposes. This is needed because -// docbook uses c++ and html5 uses cpp. -:language: {basebackend@docbook:c++:cpp} - -// This is necessary for asciidoc, but not for asciidoctor -:cpp: C++ - -== Notice - -IMPORTANT: This specification is a draft. - -Copyright (c) 2021 Intel Corporation. All rights reserved. - -NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are -trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. -used by permission by Khronos. - -== Dependencies - -This extension is written against the SYCL 2020 specification, Revision 3. - -== Status - -Draft - -This is a preview extension specification, intended to provide early access to -a feature for review and community feedback. When the feature matures, this -specification may be released as a formal extension. - -Because the interfaces defined by this specification are not final and are -subject to change they are not intended to be used by shipping software -products. - -== Version - -Revision: 3 - -== Introduction - -This extension adds functionality to convert value of single-precision -floating-point type(`float`) to `bfloat16` type and vice versa. The extension -doesn't add support for `bfloat16` type as such, instead it uses 16-bit integer -type(`uint16_t`) as a storage for `bfloat16` values. - -The purpose of conversion from float to bfloat16 is to reduce ammount of memory -required to store floating-point numbers. Computations are expected to be done with -32-bit floating-point values. - -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 -command (e.g. from `parallel_for`). - -== Feature test macro - -This extension provides a feature-test macro as described in the core SYCL -specification section 6.3.3 "Feature test macros". Therefore, an implementation -supporting this extension must predefine the macro -`SYCL_EXT_INTEL_BF16_CONVERSION` to one of the values defined in the table -below. Applications can test for the existence of this macro to determine if -the implementation supports this feature, or applications can test the macro’s - value to determine which of the extension’s APIs the implementation supports. - -[%header,cols="1,5"] -|=== -|Value |Description -|1 |Initial extension version. Base features are supported. -|=== - -== Extension to `enum class aspect` - -[source] ----- -namespace sycl { -enum class aspect { - ... - ext_intel_bf16_conversion -} -} ----- - -If a SYCL device has the `ext_intel_bf16_conversion` aspect, then it natively -supports conversion of values of `float` type to `bfloat16` and back. - -If the device doesn't have the aspect, objects of `bfloat16` class must not be -used in the device code. - -**NOTE**: The `ext_intel_bf16_conversion` aspect is not yet supported. The -`bfloat16` class is currently supported only on Xe HP GPU. - -== New `bfloat16` class - -The `bfloat16` class below provides the conversion functionality. Conversion -from `float` to `bfloat16` is done with round to nearest even(RTE) rounding -mode. - -[source] ----- -namespace sycl { -namespace ext { -namespace intel { -namespace experimental { - -class bfloat16 { - using storage_t = uint16_t; - storage_t value; - -public: - 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); - - // Convert from float to bfloat16 - bfloat16(const float &a); - bfloat16 &operator=(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-(bfloat16 &bf) { /* ... */ } - - // OP is: prefix ++, -- - friend bfloat16 &operatorOP(bfloat16 &bf) { /* ... */ } - - // OP is: postfix ++, -- - friend bfloat16 operatorOP(bfloat16 &bf, int) { /* ... */ } - - // OP is: +=, -=, *=, /= - friend bfloat16 &operatorOP(bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ } - - // OP is +, -, *, / - friend bfloat16 operatorOP(const bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ } - template - friend bfloat16 operatorOP(const bfloat16 &lhs, const T &rhs) { /* ... */ } - template - friend bfloat16 operatorOP(const T &lhs, const bfloat16 &rhs) { /* ... */ } - - // OP is ==,!=, <, >, <=, >= - friend bool operatorOP(const bfloat16 &lhs, const bfloat16 &rhs) { /* ... */ } - template - friend bool operatorOP(const bfloat16 &lhs, const T &rhs) { /* ... */ } - template - friend bool operatorOP(const T &lhs, const bfloat16 &rhs) { /* ... */ } -}; - -} // namespace experimental -} // namespace intel -} // namespace ext -} // 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 float& a);` -| Construct `bfloat16` from `float`. Converts `float` to `bfloat16`. - -| `bfloat16 &operator=(const float &a);` -| Replace the value with `a` converted 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 - 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 - 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 - 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 - 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] ----- -#include -#include - -using sycl::ext::intel::experimental::bfloat16; - -bfloat16 operator+(const bfloat16 &lhs, const bfloat16 &rhs) { - return static_cast(lhs) + static_cast(rhs); -} - -float foo(float a, float b) { - // Convert from float to bfloat16. - bfloat16 A {a}; - bfloat16 B {b}; - - // Convert A and B from bfloat16 to float, do addition on floating-pointer - // numbers, then convert the result to bfloat16 and store it in C. - bfloat16 C = A + B; - - // Return the result converted from bfloat16 to float. - return C; -} - -int main (int argc, char *argv[]) { - float data[3] = {7.0, 8.1, 0.0}; - sycl::device dev; - sycl::queue deviceQueue{dev}; - sycl::buffer buf {data, sycl::range<1> {3}}; - - if (dev.has(sycl::aspect::ext_intel_bf16_conversion)) { - deviceQueue.submit ([&] (sycl::handler& cgh) { - auto numbers = buf.get_access (cgh); - cgh.single_task ([=] () { - numbers[2] = foo(numbers[0], numbers[1]); - }); - }); - } - return 0; -} ----- - -== Issues - -None. - -== Revision History - -[cols="5,15,15,70"] -[grid="rows"] -[options="header"] -|======================================== -|Rev|Date|Author|Changes -|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 -|3|2021-08-18|Alexey Sotkin |Remove `uint16_t` constructor -|======================================== From e433fbcac8a3019510889a7f0a4db4c5f466393f Mon Sep 17 00:00:00 2001 From: "jack.kirk" Date: Thu, 31 Mar 2022 12:02:37 +0100 Subject: [PATCH 11/11] typo --- .../extensions/experimental/sycl_ext_oneapi_bfloat16.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16.asciidoc index 175219e23c47f..88b6c73b02514 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16.asciidoc @@ -57,7 +57,7 @@ floating-point type(`float`) to `bfloat16` type and vice versa. The extension doesn't add support for `bfloat16` type as such, instead it uses 16-bit integer type(`uint16_t`) as a storage for `bfloat16` values. -The purpose of conversion from float to bfloat16 is to reduce ammount of memory +The purpose of conversion from float to bfloat16 is to reduce the amount of memory required to store floating-point numbers. Computations are expected to be done with 32-bit floating-point values.