Conversation
| #ifdef __SYCL_DEVICE_ONLY__ | ||
| // New FP8 builtins | ||
| extern __DPCPP_SYCL_EXTERNAL sycl::half | ||
| __builtin_spirv_ClampConvertE4M3ToFP16INTEL(uint8_t) noexcept; |
There was a problem hiding this comment.
I was just talking with @bashbaug today, and he mentioned that there is a multi-vendor SPIR-V extension for FP8 conversions. It looks like this PR is using the proposed Intel extension. I think this is the multi-vendor extension:
https://github.khronos.org/SPIRV-Registry/extensions/EXT/SPV_EXT_float8.html
@bashbaug do you know if our driver supports the multi-vendor extension? @dklochkov-emb does the multi-vendor extension provide all the functionality we need, or would we still need something from the Intel extension?
There was a problem hiding this comment.
do you know if our driver supports the multi-vendor extension?
I believe there is some experimental support, but the extension is not formally supported currently.
It would be good to understand:
- Whether any functionality is fundamentally missing from the EXT extension. For example, the EXT extension currently does not include any stochastic rounding functionality. Is this needed, and if so, can it be provided by a layered extension?
- Any details about required support with the EXT extension. For example, which rounding modes should be supported, and which other floating-point types should be supported for conversions to/from the fp8 types?
There was a problem hiding this comment.
For example, the EXT extension currently does not include any stochastic rounding functionality. Is this needed, and if so, can it be provided by a layered extension?
Yes, we need stochastic rounding for E5M2 when converting from half, bfloat16, and float. I suppose it could be implemented by a layered extension, but it would be up to use to define such an extension.
For example, which rounding modes should be supported
For E45M3, we need only RNE with saturation.
For E5M2, we need only RNE, but we need both saturation and non-saturation.
and which other floating-point types should be supported for conversions to/from the fp8 types
I think the Intel SPIRV extension supports conversions to/from half, bfloat16, and float.
It seems like we should do one of two things:
-
Do all of the following:
- Fully implement the EXT extension.
- Define a layered SPIRV extension with the missing features.
- Change this PR to use the EXT and that new layered extension.
-
Clean up and publish the existing Intel SPIRV extension.
This PR adds FP8 types for CRI docs.
It adds the first part which includes:
This PR does not include e2e tests and checks behavior on the device due to driver and spirv translator issues - they will be added in the second part.