[SYCL] Add round-to-nearest saturating float->int8 conversions#22759
Open
zjin-lcf wants to merge 3 commits into
Open
[SYCL] Add round-to-nearest saturating float->int8 conversions#22759zjin-lcf wants to merge 3 commits into
zjin-lcf wants to merge 3 commits into
Conversation
The compare_exchange definition guard grouped __NVPTX__ with __SPIR32__ and only emitted atomics for 32-bit types, so 64-bit integer __clc_atomic_compare_exchange was declared (via atomic_decl.inc) but never defined for CUDA, causing "Unresolved extern function '_Z29__clc_atomic_compare_exchange...'" (ptxas fatal) for 64-bit CAS. Match the guard used by the sibling clc_atomic_def.inc (restrict only __SPIR32__ to 32-bit) so both CUDA (NVPTX) and HIP (AMDGPU) get 64-bit compare_exchange definitions. Co-authored-by: Cursor <cursoragent@cursor.com>
Add `sycl_ext_oneapi_saturating_conversion` helpers that round a float to
the nearest integer (round-to-nearest-even) and saturate to an 8-bit
integer, with NaN mapped to 0:
- `float_to_int8_rn(float) -> int8_t`
- `float_to_uint8_rn(float) -> uint8_t`
- `float4_to_int8x4_rn(vec<float,4>) -> int32_t`
- `float4_to_uint8x4_rn(vec<float,4>) -> uint32_t`
The packed variants produce the byte layout consumed by the
sycl_ext_oneapi_dot_accumulate `dot_acc` (dp4a) helpers, which is the
common INT8 quantization -> dp4a pipeline.
On NVPTX these lower to the single `cvt.rni.sat.{s8,u8}.f32` instruction.
Other targets (AMDGCN, host, SPIR-V) use a portable round + clamp +
convert sequence that yields bit-identical results (e.g. v_rndne_f32 +
v_max_f32/v_min_f32 + v_cvt_i32_f32 on AMDGCN). Compiler builtins are
used so the header requires no libsycl linkage.
Resolves intel#9101.
Add check_device_code codegen tests (CUDA sm_90, HIP gfx90a) and an E2E
numerical test covering saturation, ties-to-even and NaN.
Co-authored-by: Cursor <cursoragent@cursor.com>
Add the experimental extension specification for the round-to-nearest saturating float->int8 conversions. Co-authored-by: Cursor <cursoragent@cursor.com>
wenju-he
reviewed
Jul 25, 2026
Contributor
There was a problem hiding this comment.
could you also upstream this change to llvm-project?
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds round-to-nearest-even, saturating
float -> int8/uint8conversions, resolving #9101 (request for afloat_to_int8_rn-style helper matching CUDA'scvt.rni.sat.s8.f32).New helpers in
sycl::ext::oneapi::experimental:int8_t float_to_int8_rn(float)uint8_t float_to_uint8_rn(float)int32_t float4_to_int8x4_rn(vec<float,4>)uint32_t float4_to_uint8x4_rn(vec<float,4>)The packed 4-wide variants produce the byte layout consumed by the
sycl_ext_oneapi_dot_accumulatedot_acc/dp4a helpers, i.e. the typical INT8 quantization -> dp4a pipeline.Semantics (identical on all backends): round-to-nearest-even, saturate to the destination range,
NaN -> 0.Lowering:
cvt.rni.sat.s8.f32/cvt.rni.sat.u8.f32.f32->sat.s8instruction exists on AMDGCN); lowers tov_rndne_f32+v_max_f32/v_min_f32+v_cvt_i32_f32. Compiler builtins are used so the header needs no libsycl linkage.Test plan
check_device_code/cuda/saturating_conversion.cpp— FileCheckcvt.rni.sat.{s8,u8}.f32(sm_90).check_device_code/hip/saturating_conversion.cpp— FileCheck the round+clamp+convert IR (gfx90a); ISA verified to containv_rndne_f32/v_max_f32/v_min_f32/v_cvt_i32_f32.test-e2e/SaturatingConversion/saturating_conversion.cpp— numerical validation vs. a host reference over saturation, ties-to-even and NaN. Passes on:Notes
A formal extension specification document can be added if maintainers would like this promoted beyond the experimental namespace.