[SYCL] Implement `sycl_ext_oneapi_fp8` for CRI by dklochkov-emb · Pull Request #21568 · intel/llvm

dklochkov-emb · 2026-03-19T13:48:19Z

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

Initial implementation
Unit tests which check expected behavior and values
Unit tests which check the fact of builtin calls.
E2E tests of all 3 data types: fp8_e4m3, fp8_e5m2 and fp8_e8m0

…8-new-arch

gmlueck · 2026-03-20T15:13:09Z

+#ifdef __SYCL_DEVICE_ONLY__
+// New FP8 builtins
+extern __DPCPP_SYCL_EXTERNAL sycl::half
+__builtin_spirv_ClampConvertE4M3ToFP16INTEL(uint8_t) noexcept;


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?

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?

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.

I believe we already have everything ready in terms of SPIR-V specs and support -- SPV_INTEL_fp_conversions was published in December 2025 as a layered extension on top of SPV_EXT_float8, and it covers all the functionality you mention.

Specifically:

RNE without saturation (E5M2): OpFConvert from SPV_EXT_float8

RNE with saturation (E4M3 and E5M2): OpClampConvertFToFINTEL from SPV_INTEL_fp_conversions

Stochastic rounding (E5M2, with and without saturation): OpStochasticRoundFToFINTEL/OpClampStochasticRoundFToFINTEL from SPV_INTEL_fp_conversions

Source types (half, bfloat16, float): supported in both extensions

EXT does not support Stochastic rounding, @gmlueck @bashbaug , can it be used INTEL spirv in this case?

I agree that the EXT extension does not support stochastic rounding. It should be OK to use OpStochasticRoundFToFINTEL and OpClampStochasticRoundFToFINTEL for this instead, though.

SPV_INTEL_fp_conversions was published in December 2025 as a layered extension on top of SPV_EXT_float8, and it covers all the functionality you mention.

This extension is only "approved for public preview" at present. Since we plan to use it in SYCL, it should be promoted to "approved for publication", and it should be published on the Khronos website.

RNE with saturation (E4M3 and E5M2): OpClampConvertFToFINTEL from SPV_INTEL_fp_conversions

It seems like this functionality is also provided by the SPV_EXT_float8 extension via the SaturatedToLargestFloat8NormalConversionEXT decoration. If we want SPV_INTEL_fp_conversions to contain only additional functionality, should we remove OpClampConvertFToFINTEL? CC: @bashbaug

It also seems like there is some functionality required by SYCL that is missing from both SPIRV extensions. The SYCL extension requires saturation (clamped) conversion when converting to integer. The INTEL extension provides OpClampConvertFToSINTEL, but that only works when converting to a signed integer. For some reason, there is no clamped equivalent to OpConvertFToU. Was this an oversight in the SPIRV extension?

I was trying to remember why I specified the SYCL extension to require saturation (clamped) conversion for integers. C++ says that FP to integer conversions are undefined when the value cannot be represented after rounding. I think I required saturation because this is how the CUDA API works. If you scroll way down to "Integer notes", you can see that the PTX cvt instruction guarantees saturation when converting to an integer value.

Updated to use EXT where it is possible

This reverts commit a55e275.

dklochkov-emb · 2026-03-31T08:17:09Z

@intel/llvm-reviewers-runtime please, review

Copilot

Pull request overview

This PR introduces initial host-side implementations of OneAPI experimental FP8 types (E4M3, E5M2, E8M0) and adds unit tests validating conversions/encodings plus tests that verify the expected SPIR-V builtin call paths (via mocks) for device-only code paths.

Changes:

Added sycl::ext::oneapi::experimental FP8 type definitions and conversion logic in a new public header.
Added unit tests for value/encoding behavior for fp8 E4M3, E5M2, and E8M0 on CPU.
Added mocked builtin implementations and tests to confirm which builtins are invoked from constructors/conversions.

Reviewed changes

Copilot reviewed 8 out of 8 changed files in this pull request and generated 8 comments.

Show a summary per file

File	Description
`sycl/include/sycl/ext/oneapi/experimental/float_8bit/types.hpp`	Adds the FP8 type implementations, conversion helpers, and builtin hooks.
`sycl/unittests/Extensions/fp8/fp8_e4m3.cpp`	Adds host unit tests for E4M3 encoding/conversion behavior.
`sycl/unittests/Extensions/fp8/fp8_e5m2.cpp`	Adds host unit tests for E5M2 encoding/conversion behavior.
`sycl/unittests/Extensions/fp8/fp8_e8m0.cpp`	Adds host unit tests for E8M0 encoding/conversion behavior.
`sycl/unittests/Extensions/fp8/builtin_mocks.hpp`	Provides mock builtin symbols and counters to validate builtin call paths.
`sycl/unittests/Extensions/fp8/builtin_call_tests.cpp`	Adds tests asserting the expected builtin calls occur (via the mocks).
`sycl/unittests/Extensions/fp8/CMakeLists.txt`	Registers the new FP8 unit test target and sources.
`sycl/unittests/Extensions/CMakeLists.txt`	Adds the new `fp8` unit test subdirectory to the Extensions suite.

…8-new-arch

gmlueck

Here are my comments so far. I haven't looked at E8M0 yet, or at any of the CPU conversions.

Co-authored-by: Copilot <copilot@github.com>

…8-new-arch

Co-authored-by: Copilot <copilot@github.com>

maksimsab

SYCLPostLink part LGTM.

Co-authored-by: Copilot <copilot@github.com>

dklochkov-emb · 2026-05-19T09:39:33Z

E2E tests were added but it is required driver version with bug fixes to pass on CI. Locally all tests pass

Co-authored-by: Copilot <copilot@github.com>

dklochkov-emb · 2026-05-25T12:13:30Z

@gmlueck Please, review it

dm-vodopyanov · 2026-05-26T12:57:49Z

+  explicit fp8_e8m0_x(half const (&in)[N], rounding r = rounding::upward) {
+    CheckConstraints(r);
+    for (size_t i = 0; i < N; ++i)
+      vals[i] = detail::ConvertFloatToE8M0_CPU(in[i], r, saturation::finite);
+  }
+
+  explicit fp8_e8m0_x(bfloat16 const (&in)[N], rounding r = rounding::upward) {
+    CheckConstraints(r);
+    for (size_t i = 0; i < N; ++i)
+      vals[i] = detail::ConvertFloatToE8M0_CPU(in[i], r, saturation::finite);
+  }
+
+  explicit fp8_e8m0_x(float const (&in)[N], rounding r = rounding::upward) {
+    CheckConstraints(r);
+    for (size_t i = 0; i < N; ++i)
+      vals[i] = detail::ConvertFloatToE8M0_CPU(in[i], r, saturation::finite);
+  }
+
+  explicit fp8_e8m0_x(const marray<half, N> &in,
+                      rounding r = rounding::upward) {
+    CheckConstraints(r);
+    for (size_t i = 0; i < N; ++i)
+      vals[i] = detail::ConvertFloatToE8M0_CPU(in[i], r, saturation::finite);
+  }
+
+  explicit fp8_e8m0_x(const marray<bfloat16, N> &in,
+                      rounding r = rounding::upward) {
+    CheckConstraints(r);
+    for (size_t i = 0; i < N; ++i)
+      vals[i] = detail::ConvertFloatToE8M0_CPU(in[i], r, saturation::finite);
+  }
+
+  explicit fp8_e8m0_x(const marray<float, N> &in,
+                      rounding r = rounding::upward) {
+    CheckConstraints(r);
+    for (size_t i = 0; i < N; ++i)
+      vals[i] = detail::ConvertFloatToE8M0_CPU(in[i], r, saturation::finite);
+  }


Maybe use templates with #define to avoid code duplication in the entire file?

gmlueck

I didn't finish reviewing this yet, but I wanted to send the comments I have so far. So far, I have just reviewed the E4M3 and E5M2 device-side code.

gmlueck · 2026-05-26T19:37:37Z

+#ifdef __SYCL_DEVICE_ONLY__
+    _Float16 v{0};
+    if constexpr (std::is_same_v<std::decay_t<T>, sycl::half>)
+      v = static_cast<_Float16>(static_cast<float>(h));


I don't understand this line. Isn't half the same as _Float16? Why do we convert from half to float and then back to _Float16?

gmlueck · 2026-05-26T19:44:35Z

+    if constexpr (std::is_same_v<std::decay_t<T>, sycl::half>)
+      v = static_cast<_Float16>(static_cast<float>(h));
+    else
+      v = static_cast<_Float16>(h);


When T is float, doesn't this have the same two-step conversion problem that is described here: https://github.com/intel-tools/intel-xpu-backend-for-triton/issues/847 ?

I'm not sure if there are also similar problems when T is an integer type. Did you consider this and decide it is safe?

gmlueck · 2026-05-26T19:53:47Z

+    if constexpr (((std::is_same_v<std::decay_t<Types>, bfloat16>) && ...)) {
+      const bfloat16 in[N] = {static_cast<bfloat16>(v)...};
+      for (size_t i = 0; i < N; ++i)
+        vals[i] = ConvertBF16ToFP8(in[i]);


The SPIR-V instructions __builtin_spirv_ClampConvertFP16ToE4M3INTEL and __builtin_spirv_ClampConvertBF16ToE4M3INTEL allow the input to be a vector. In the N == 2 case, I wonder if we are losing some performance by calling the instruction twice (each with a scalar input) rather than calling it once (with a vector of two elements).

gmlueck · 2026-05-26T20:09:21Z

+  explicit operator sycl::marray<sycl::half, N>() const {
+    sycl::marray<sycl::half, N> ret;
+    for (size_t i = 0; i < N; ++i)
+      ret[i] = ConvertFromFP8<sycl::half>(vals[i]);


Same comment here about performance. I wonder if we are losing some performance by calling this in a loop rather than using the vector versions of the conversion instructions.

gmlueck · 2026-05-26T20:22:41Z

+      }
+      current_seed = *seed.pseed;
+    }
+#endif


We should either assert or throw an exception on host.

gmlueck · 2026-05-26T20:26:00Z

+#ifdef __SYCL_DEVICE_ONLY__
+    uint32_t current_seed = *seed.pseed;
+    for (size_t i = 0; i < N; ++i) {
+      const _Float16 v = static_cast<_Float16>(static_cast<float>(in[i]));


Same question as earlier, why convert from half to float and back to _Float16?

gmlueck · 2026-05-27T14:24:28Z

+  // Intentionally public to allow access to the raw values.
+  uint8_t vals[N];
+};
+


I think the deduction guides are missing for all three of these classes. The spec calls for this deduction guide for the fp8_e4m3_x class, for example:

// Deduction guide available only when the size of the pack is greater than zero. template<typename... Ts> fp8_e4m3_x(Ts...) -> fp8_e4m3_x<sizeof...(Ts)>;

gmlueck

All these comments are about the SPIR-V intrinsics. We should use the SPV_EXT_float8 extension as much as possible and only use SPV_INTEL_fp_conversions for the things it does not contain. It's not clear to me which SPIR-V instructions are being emitted from these intrinsics, but some of the names make it sound like you may be emitting SPV_INTEL_fp_conversions instead of SPV_EXT_float8.

I've listed the SPIR-V instructions I think we should use for each of the intrinsics below.

Tagging @bashbaug here also to make sure I got this right. :-)

gmlueck · 2026-05-27T15:15:11Z

+// FP8 builtins
+
+extern __DPCPP_SYCL_EXTERNAL uint8_t
+__builtin_spirv_ClampConvertFP16ToE4M3INTEL(_Float16) noexcept;


I think this should be using OpFConvert with Float8E4M3EXT as the destination FP Encoding.

I see that the name of this builtin ends with "INTEL", so I presume it is not emitting OpFConvert with Float8E4M3EXT now?

gmlueck · 2026-05-27T15:24:54Z

+__builtin_spirv_ClampConvertFP16ToE4M3INTEL(_Float16) noexcept;
+
+extern __DPCPP_SYCL_EXTERNAL _Float16
+__builtin_spirv_ConvertE4M3ToFP16EXT(char) noexcept;


And this should be using OpFconvert with Float8E4M3EXT as the source FP Encoding.

I see that the name of this builtin ends with "EXT". Does that mean it is already emitting this SPIR-V instruction?

gmlueck · 2026-05-27T15:26:06Z

+__builtin_spirv_ConvertE4M3ToFP16EXT(char) noexcept;
+
+extern __DPCPP_SYCL_EXTERNAL uint8_t
+__builtin_spirv_ClampConvertBF16ToE4M3INTEL(__bf16) noexcept;


This should use OpFconvert with Float8E4M3EXT as the destination FP Encoding and decorated with SaturatedToLargestFloat8NormalConversionEXT.

gmlueck · 2026-05-27T15:27:36Z

+extern __DPCPP_SYCL_EXTERNAL uint8_t
+__builtin_spirv_ClampConvertBF16ToE4M3INTEL(__bf16) noexcept;
+extern __DPCPP_SYCL_EXTERNAL __bf16
+__builtin_spirv_ConvertE4M3ToBF16EXT(uint8_t) noexcept;


This should use OpFConvert with Float8E4M3EXT as the source FP Encoding and BFloat16KHR as the destination FP Encoding.

gmlueck · 2026-05-27T15:28:36Z

+extern __DPCPP_SYCL_EXTERNAL __bf16
+__builtin_spirv_ConvertE4M3ToBF16EXT(uint8_t) noexcept;
+extern __DPCPP_SYCL_EXTERNAL uint8_t
+__builtin_spirv_ClampConvertFP16ToE5M2INTEL(_Float16) noexcept;


This should use OpFConvert with Float8E5M2EXT as the destination FP Encoding and the SaturatedToLargestFloat8NormalConversionEXT decoration.

gmlueck · 2026-05-27T15:32:09Z

+extern __DPCPP_SYCL_EXTERNAL _Float16
+__builtin_spirv_ConvertE5M2ToFP16EXT(uint8_t) noexcept;
+extern __DPCPP_SYCL_EXTERNAL uint8_t
+__builtin_spirv_ClampConvertBF16ToE5M2INTEL(__bf16) noexcept;


This should use OpFConvert with BFloat16KHR as the source FP Encoding and Float8E5M2EXT as the destination FP Encoding and decorated with SaturatedToLargestFloat8NormalConversionEXT.

gmlueck · 2026-05-27T15:33:26Z

+extern __DPCPP_SYCL_EXTERNAL uint8_t
+__builtin_spirv_ClampConvertBF16ToE5M2INTEL(__bf16) noexcept;
+extern __DPCPP_SYCL_EXTERNAL uint8_t
+__builtin_spirv_ConvertBF16ToE5M2EXT(__bf16) noexcept;


This should use OpFConvert with BFloat16KHR as the source FP Encoding and Float8E5M2EXT as the destination FP Encoding.

gmlueck · 2026-05-27T15:34:52Z

+extern __DPCPP_SYCL_EXTERNAL uint8_t
+__builtin_spirv_ConvertBF16ToE5M2EXT(__bf16) noexcept;
+extern __DPCPP_SYCL_EXTERNAL __bf16
+__builtin_spirv_ConvertE5M2ToBF16EXT(uint8_t) noexcept;


This should use OpFConvert with Float8E5M2EXT as the source FP Encoding and BFloat16KHR as the destination FP Encoding.

gmlueck · 2026-05-27T15:43:45Z

+__builtin_spirv_ConvertE5M2ToBF16EXT(uint8_t) noexcept;
+extern __DPCPP_SYCL_EXTERNAL uint8_t
+__builtin_spirv_ClampStochasticRoundFP16ToE5M2INTEL(_Float16, uint32_t,
+                                                    uint32_t *) noexcept;


This should use OpStochasticRoundFToFINTEL with the SaturatedToLargestFloat8NormalConversionEXT decoration.

gmlueck · 2026-05-27T15:46:41Z

+                                               uint32_t *) noexcept;
+extern __DPCPP_SYCL_EXTERNAL uint8_t
+__builtin_spirv_ClampStochasticRoundBF16ToE5M2INTEL(__bf16, uint32_t,
+                                                    uint32_t *) noexcept;


This should use OpStochasticRoundFToFINTEL with BFloat16KHR as the source FP Encoding and the SaturatedToLargestFloat8NormalConversionEXT decoration.

gmlueck · 2026-05-27T17:02:55Z

+#include <stdexcept>
+#include <type_traits>
+
+#define SYCL_EXT_ONEAPI_FP8 1


I think we normally define these macros in "sycl/source/feature_test.hpp.in".

dklochkov-emb added 7 commits February 18, 2026 09:28

[SYCL] add new fp8 data types and unit tests

3a25ebf

[SYCL] update fp8 to check constraints

9490498

[SYCL] apply new updates from docs and e5m3 data type

a24ac3b

[SYCL] remove extra types

2573ebe

[SYCL][FP8] implement stochastic rounding

bb0cc94

[SYCL] update fp8 implemetation

6e0a2e4

Merge remote-tracking branch 'upstream/sycl' into sycl-ext-one-api-fp…

1715062

…8-new-arch

dklochkov-emb requested a review from a team as a code owner March 19, 2026 13:48

dklochkov-emb requested a review from sergey-semenov March 19, 2026 13:48

dklochkov-emb added 2 commits March 19, 2026 14:59

[SYCL] fix formatting

6d03f08

[SYCL] do not use extra rounding modes

c938642

gmlueck reviewed Mar 20, 2026

View reviewed changes

dklochkov-emb added 4 commits March 24, 2026 13:41

[SYCL][FP8] use saturation

a55e275

Revert "[SYCL][FP8] use saturation"

28565bc

This reverts commit a55e275.

[SYCL] update list of builtins used in fp8 types

9712f34

[SYCL] add more tests of builtin calls

046affd

uditagarwal97 requested a review from Copilot March 31, 2026 15:06

Copilot started reviewing on behalf of uditagarwal97 March 31, 2026 15:08 View session

Copilot AI reviewed Mar 31, 2026

View reviewed changes

dklochkov-emb added 8 commits April 1, 2026 11:50

[SYCL] fix PR issues

a697eb8

[SYCL] do not use extra check for e8m0

1f0808a

[SYCL] fix formatting

ef5f670

Merge remote-tracking branch 'upstream/sycl' into sycl-ext-one-api-fp…

865d909

…8-new-arch

[SYCL] remove unused variable

8d9cc9f

[SYCL] remove unused variable

e5fd6c4

Merge remote-tracking branch 'upstream/sycl' into sycl-ext-one-api-fp…

f9a044f

…8-new-arch

[SYCL] fix formatting

ae426a9

gmlueck reviewed Apr 6, 2026

View reviewed changes

[SYCL] do not construct fp8 with mixture of parameters in pack

5b6da23

dklochkov-emb and others added 7 commits May 12, 2026 18:14

[SYCL] remove doubles and add more tests

9caabe8

Co-authored-by: Copilot <copilot@github.com>

[SYCL][TESTS] FP8: fix builtin mock tests

113be77

Co-authored-by: Copilot <copilot@github.com>

Merge remote-tracking branch 'upstream/sycl' into sycl-ext-one-api-fp…

7272e02

…8-new-arch

[SYCL][TESTE2E] make coverage of fp8 data types about 90

9bfe2b7

[SYCL][E2E] add more FP8 e2e tests

4b44878

Merge remote-tracking branch 'upstream/sycl' into sycl-ext-one-api-fp…

2aeaba0

…8-new-arch

[SYCL] do not set port for communication with sumilator

0bbf809

dklochkov-emb requested a review from a team as a code owner May 18, 2026 08:20

dklochkov-emb changed the title ~~FP8 types for CRI PART 1~~ FP8 types for CRI May 18, 2026

dklochkov-emb and others added 4 commits May 18, 2026 10:25

[SYCL][E2E] add requirement of cri gpu

30cf7c6

Co-authored-by: Copilot <copilot@github.com>

[SYCL] fix formatting

655e8ad

[SYCL] add FP8 feature macro

f338a2f

Co-authored-by: Copilot <copilot@github.com>

[SYCL] fix formating

e9eb83f

dm-vodopyanov changed the title ~~FP8 types for CRI~~ [SYCL] Implement sycl_ext_oneapi_fp8 for CRI May 18, 2026

maksimsab requested changes May 18, 2026

View reviewed changes

Comment thread llvm/lib/SYCLPostLink/ModuleSplitter.cpp

[SYCL][TEST] update test of post link drop known builtins

0ed3ca2

Co-authored-by: Copilot <copilot@github.com>

dklochkov-emb requested a review from maksimsab May 18, 2026 14:10

maksimsab approved these changes May 18, 2026

View reviewed changes

[SYCL][E2E] add spirv translator pattern to build command

f8f8952

Co-authored-by: Copilot <copilot@github.com>

dklochkov-emb requested a review from gmlueck May 19, 2026 09:39

dklochkov-emb and others added 2 commits May 19, 2026 14:30

[SYCL][E2E] make tests XFAIL until driver is installed on CI machines

c82aa03

Co-authored-by: Copilot <copilot@github.com>

[SYCL][TESTE2E] fix formatting

880c725

[SYCL] fix build issues on win

b0f212c

dm-vodopyanov reviewed May 26, 2026

View reviewed changes

gmlueck reviewed May 26, 2026

View reviewed changes

gmlueck reviewed May 27, 2026

View reviewed changes

Conversation

dklochkov-emb commented Mar 19, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

vmaksimo Mar 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

dklochkov-emb commented Mar 31, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

gmlueck left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

maksimsab left a comment

Choose a reason for hiding this comment

Uh oh!

dklochkov-emb commented May 19, 2026

Uh oh!

dklochkov-emb commented May 25, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

gmlueck left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

gmlueck left a comment

Choose a reason for hiding this comment

dklochkov-emb commented Mar 19, 2026 •

edited

Loading

vmaksimo Mar 24, 2026 •

edited

Loading

dklochkov-emb commented Mar 31, 2026 •

edited

Loading