[SYCL] Add sycl-enable-bfloat16-conversion compilation flag

[MrSidims]

This flag defines SYCL_BF16_CONVERSION_IS_SUPPORTED macro that
can be used to specify whether BF16 conversion feature is supported.
This macro is being defined for both host and device compilation.

Signed-off-by: Dmitry Sidorov [email protected]

[MrSidims]

Created instead of #4266

[AaronBallman]

This flag doesn't currently have any semantics beyond defining a macro; shouldn't there also be codegen changes as part of this?

[MrSidims]

This flag doesn't currently have any semantics beyond defining a macro; shouldn't there also be codegen changes as part of this?

No, it's just about definition of the macro.
A little backstory. We are adding a new experimental feature (#4213) that enables conversion from float to bfloat16 and vice versa. We want to give a user a possibility to pick which version of the conversion function to pick: DPCPP compiler's or some backup implementation (implemented by user), so on the user's end it would be:

uint16_t convert_to_bf16(float a) {
#ifdef __SYCL_ENABLE_BF16_CONVERSION__
return cl::sycl::ext::intel::experimental::bfloat16::from_float(Val);
#else
return users_from_float(Val);
#endif // __SYCL_ENABLE_BF16_CONVERSION__
}

At some point of time we may add a fallback implementation in SYCL headers as well.
We are introducing this compilation flag because:
0. this way this macro will be documented by the compiler (and may be we will re-use this macro in SYCL headers);

1. there is no low-level runtime specification yet to have SYCL2020 aspect mapping;
2. a better way would be re-use device optional feature extension (https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/DeviceIf/device_if.asciidoc) but it's not yet implemented.

[AaronBallman]

This flag doesn't currently have any semantics beyond defining a macro; shouldn't there also be codegen changes as part of this?

No, it's just about definition of the macro.
A little backstory. We are adding a new experimental feature (#4213) that enables conversion from float to bfloat16 and vice versa. We want to give a user a possibility to pick which version of the conversion function to pick: DPCPP compiler's or some backup implementation (implemented by user), so on the user's end it would be:

uint16_t convert_to_bf16(float a) {
#ifdef __SYCL_ENABLE_BF16_CONVERSION__
return cl::sycl::ext::intel::experimental::bfloat16::from_float(Val);
#else
return users_from_float(Val);
#endif // __SYCL_ENABLE_BF16_CONVERSION__
}

At some point of time we may add a fallback implementation in SYCL headers as well.
We are introducing this compilation flag because:
0. this way this macro will be documented by the compiler (and may be we will re-use this macro in SYCL headers);

1. there is no low-level runtime specification yet to have SYCL2020 aspect mapping;
2. a better way would be re-use device optional feature extension (https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/DeviceIf/device_if.asciidoc) but it's not yet implemented.

Thanks for the background! Is the plan then to deprecate this option once we implement the optional device feature extension? If so, is there a pressing need for landing this as official functionality (as opposed to experimenting with it in a branch or something until the general feature lands)?

[MrSidims]

Thanks for the background! Is the plan then to deprecate this option once we implement the optional device feature extension?

The plan is to set SYCLEnableBF16Conversion's default value to "true" when the feature matures leaving a way to disable it if necessary with -fno option.

If so, is there a pressing need for landing this as official functionality (as opposed to experimenting with it in a branch or something until the general feature lands)?

I'd say, not a 'pressing need', but it would be really appreciated by our customers.

[MrSidims]

I would also like to hear from @AlexeySachkov as he has a better knowledge of optional device's features status/design. Do you see any drawbacks of introducing this option?

[AaronBallman]

Thanks for the background! Is the plan then to deprecate this option once we implement the optional device feature extension?

The plan is to set SYCLEnableBF16Conversion's default value to "true" when the feature matures leaving a way to disable it if necessary with -fno option.

If so, is there a pressing need for landing this as official functionality (as opposed to experimenting with it in a branch or something until the general feature lands)?

I'd say, not a 'pressing need', but it would be really appreciated by our customers.

I won't block this change, but I would say that my general preference is to never add a feature that we already know will be replaced by another approach in the foreseeable future.

[AaronBallman]

LGTM aside from concerns raised about timing.

[hchilama]

Driver LGTM

[AlexeySachkov]

I would also like to hear from @AlexeySachkov as he has a better knowledge of optional device's features status/design. Do you see any drawbacks of introducing this option?

The existence of the option and support burden it brings to us :) I think I'm with @AaronBallman here. BTW, why can't we just document a particular macro which our customers could pass/pass with particular value/explicitly undefine/pass -DNOMACRO so the header file content could still be controlled by users, but we will avoid changes to the compiler?

[MrSidims]

I would also like to hear from @AlexeySachkov as he has a better knowledge of optional device's features status/design. Do you see any drawbacks of introducing this option?

The existence of the option and support burden it brings to us :) I think I'm with @AaronBallman here. BTW, why can't we just document a particular macro which our customers could pass/pass with particular value/explicitly undefine/pass -DNOMACRO so the header file content could still be controlled by users, but we will avoid changes to the compiler?

Okay, I'll check internally whether it's acceptable. @AlexeySachkov or @AaronBallman could you please 'request changes' so the PR couldn't be merged by incident?

[AaronBallman]

Requesting changes to avoid an accidental merge while we think on the timing.

[AlexeySachkov]

Explicitly requesting changes to wait until the discussion about command line option vs a macro is settled down

[MrSidims]

No need for the patch