bfloat16 data type support in SYCL

[huangzhiyuan]

As title, I know SYCL implementation defines half data type in cl::sycl space, my question is bfloat16 is whether supported like cl::sycl::bfloat16 ?

[keryell]

It is not supported for now, but someone can come with an extension for it.
More generally it would be nice to have quite more data types, especially when the hardware allows some great acceleration with them (arbitrary precision integer and fixed-point on FPGA...).

[fraggamuffin]

dpc++ has this extension
aiming SYCL Next

[fraggamuffin]

https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16.asciidoc

[Pennycook]

SYCL-Next will allow implementations to use versions of C++ other than C++17, and C++23 defines std::bfloat16_t.

Given that we've moved away from pre-adopting features from C++ in favor of just using features from C++ directly, is there still anything for us to do here? My inclination here would be to close this issue, since C++ beat us to it.

What do you think, @gmlueck?

[bader]

It's not clear what exactly "using features from C++ directly" means here.
I think we can expect that std::bfloat16_t can be used as an argument of arithmetic operations +, -, /, *. Probably C++ conversions from/to other C++ fundamental data types (int, float, etc.).
Can we use it with vec<> class? Does it mean that math functions must support std::bfloat16_t? If so, which functions?
FYI: DPC++ extension supports limited set of math functions - https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16_math_functions.asciidoc.

[keryell]

Is there an implementation where this just works right now? @illuhad

[illuhad]

@keryell I imagine that in the AdaptiveCpp SSCP compiler, it will happily generate bfloat16 types in LLVM IR in device code, and, if the LLVM backend supports it (how widely available is support in LLVM GPU backends currently?), it might just work for basic arithmetic instructions.

However it's entirely untested, I've never tried it, and we don't implement any math builtins for bf16.

[Pennycook]

It's not clear what exactly "using features from C++ directly" means here. I think we can expect that std::bfloat16_t can be used as an argument of arithmetic operations +, -, /, *. Probably C++ conversions from/to other C++ fundamental data types (int, float, etc.).

All I meant was that it seems like we shouldn't introduce a sycl::bfloat16 type, because we can just use std::bfloat16_t instead. (We should also deprecate sycl::half in favor of std::float16_t).

Can we use it with vec<> class? Does it mean that math functions must support std::bfloat16_t? If so, which functions? FYI: DPC++ extension supports limited set of math functions - https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16_math_functions.asciidoc.

You're right, we should answer these questions. But my preference would be to not make a special case for std::bfloat16_t, and instead make SYCL-Next changes to the effect of:

• All fundamental types are supported in device code.
• The availability of specific fundamental types (e.g. std::bfloat16_t) depends on the C++ version.
• sycl::vec supports all fundamental types.
• The SYCL math functions support all fundamental types.

[illuhad]

You're right, we should answer these questions. But my preference would be to not make a special case for std::bfloat16_t, and instead make SYCL-Next changes to the effect of:

Uh.. I suspect this would make it very hard to support SYCL with recent C++ versions on devices where bfloat16 is not natively supported. We'd have to implement emulations or automatic widening to float or so.

[Pennycook]

Uh.. I suspect this would make it very hard to support SYCL with recent C++ versions on devices where bfloat16 is not natively supported. We'd have to implement emulations or automatic widening to float or so.

Sorry, I could have been clearer here.

Remember that "supported in device code" is not the same thing as "supported by all devices". All I'm saying is that we should handle std::bfloat16_t the way we treat other fundamental types (e.g., double): the math functions, vec, etc all support the type, and the specification defines the necessary overloads, but we can add an aspect to check for compatible devices.

So, std::bfloat16_t would work iff somebody was using a C++23 compiler and the device had a bf16 aspect.

[illuhad]

So, std::bfloat16_t would work iff somebody was using a C++23 compiler and the device had a bf16 aspect.

Ah. I missed that the idea is to also introduce an aspect. So we'd mirror what we do for half currently (plus the C++ version constraint). That seems fine to me.