Skip to content

FP8 types for CRI PART 1#21568

Open
dklochkov-emb wants to merge 9 commits intointel:syclfrom
dklochkov-emb:sycl-ext-one-api-fp8-new-arch
Open

FP8 types for CRI PART 1#21568
dklochkov-emb wants to merge 9 commits intointel:syclfrom
dklochkov-emb:sycl-ext-one-api-fp8-new-arch

Conversation

@dklochkov-emb
Copy link
Contributor

@dklochkov-emb dklochkov-emb commented Mar 19, 2026

This PR adds FP8 types for CRI docs.
It adds the first part which includes:

  1. Initial implementation
  2. Unit tests which check expected behavior and values
  3. Unit tests which check the fact of builtin calls.

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.

@dklochkov-emb dklochkov-emb requested a review from a team as a code owner March 19, 2026 13:48
#ifdef __SYCL_DEVICE_ONLY__
// New FP8 builtins
extern __DPCPP_SYCL_EXTERNAL sycl::half
__builtin_spirv_ClampConvertE4M3ToFP16INTEL(uint8_t) noexcept;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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:

  1. 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?
  2. 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?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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:

  1. 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.
  2. Clean up and publish the existing Intel SPIRV extension.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants