[SYCL][ESIMD] Fix compilation break occurring when bfloat16 constructor is used in a kernel

[fineg74]

No description provided.

[sarnex]

Sorry for the late review, I missed this somehow.

Can you bring me up to speed on the root cause and the solution? I looked at the internal bug tracker but didn't totally understand how it's related to this PR.

My understanding is that all __devicelib_ConvertFToBF16INTEL does is call __spirv_ConvertFToBF16INTEL, at least that's what this code seems to be doing.

Do we not link against whatever sycl device library contains the implementation for __devicelib_ConvertFToBF16INTEL so we never see the defn or something?

Thanks

[fineg74]

My understanding of the problem is:
BF16 uses __devicelib_ConvertFToBF16INTEL and __devicelib_ConvertBF16ToFINTEL rather than spirv extensions for some reason. The issue is that these functions are not implemented on GPU backend while spirv extensions are implemented there. I also found implementation of these functions at llvm/libdevice/bfloat16_wrapper.cpp (wrappers from _devicelib* functions to their appropriate spirv extensions) but they are covered with SPIR define which is not set at the time of backend compilation. I am not sure why it was done this way and why. So the basic idea of the fix is to make sycl-post-link replace calls of _devicelib* functions with spirv extensions. It is slightly more complicated then usual due to 2 reasons:

1. _devicelib* functions are declared as extern "C" which means it is not mangled which causes sycl-post-link to crash as it tries to demangle the function names (it assumes the function names are always mangled). So there is code that creates a name mangling before processing to make sycl-post-link happy.
2. _devicelib* functions receive their parameters as references (not sure why as it doesn't make a lot of sense to use a reference to pass a const scalar as a parameter) while spirv extensions get their parameters by value which requires insertion of load instruction for function parameters

[sarnex]

Thanks for the explanation.

I investigated it a bit more and it seems __devicelib_ConvertBF16ToFINTEL is never translated to __spirv_ConvertFToBF16INTEL even in working non-ESIMD cases, even the GPU VISA does a function call for it, so I guess it's defined somewhere somehow.

It seems like we need to handle this on our side, so this solution makes sense, will review it in depth soon.

[sarnex]

this looks good, thanks. i dont understand why the bf16 feature was implemented this way but this looks to be a good way to deal with it.

[taozha2]

verify failed with oneapi 2023.2.0.20230622, still meet compiler issue like below:
In file included from /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/algorithm:60:
In file included from /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/utility:70:
In file included from /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/bits/stl_pair.h:59:
In file included from /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/bits/move.h:57:
/usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/type_traits:2585:44: error: no type named 'type' in 'std::enable_if'; 'enable_if' cannot be used to disable this declaration
using enable_if_t = typename enable_if<_Cond, _Tp>::type;
^~~~~
/opt/intel/oneapi/compiler/2023.2.0/linux/bin-llvm/../include/sycl/ext/intel/esimd/simd.hpp:57:23: note: in instantiation of template type alias 'enable_if_t' requested here
std::enable_if_t<detail::is_valid_simd_elem_type_v>> {
^
/home/zt/workspace/xetla_validation/libraries.gpu.xetla/build/tests/bug_track/compiler/convert_0_to_bf16.cpp:35:45: note: in instantiation of template class 'sycl::_V1::ext::intel::esimd::simd<gpu::xetla::bf16, 32>' requested here
__ESIMD_NS::simd<data_type, 32> data = data_type(0);