[SYCL] float_to_int8_rn function

[jinz2014]

I have a question. Is there a function for rounding to nearest integer with saturation in SYCL ? Thanks.

/*
 * Copyright (c) 2019-2023, NVIDIA CORPORATION.  All rights reserved.
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 *     http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */

static inline __device__ char float_to_int8_rn(float x)
{
  uint dst;
  asm volatile("cvt.rni.sat.s8.f32 %0, %1;" : "=r"(dst) : "f"(x));
  return reinterpret_cast<const char &>(dst);
}

[AlexeySachkov]

Hi @jinz2014,

I have a question. Is there a function for rounding to nearest integer with saturation in SYCL ?

I'm looking into the spec, but I don't think that there is a direct function like this in SYCL. My current findings:

4.17.5. Math functions:

genfloat rint(genfloat x)

Round to integral value (using round to nearest even rounding mode) in floating-point format. Refer to section 7.1 of the OpenCL 1.2 specification document for description of rounding modes.

Uses right rounding mode, but no saturation. Returns float, not int or char

4.14.2.1. Vec interface.

template <typename ConvertT,
          rounding_mode RoundingMode = rounding_mode::automatic>
vec<ConvertT, NumElements> convert() const

Converts this SYCL vec to a SYCL vec of a different element type specified by ConvertT using the rounding mode specified by RoundingMode. The new SYCL vec type must have the same number of elements as this SYCL vec. The different rounding modes are described in Table 147.

This should also work, but looks more hacky, I think: you can have a single-element vector and use the method to convert values to another type. Possible to specify desired rounding mode and have char return type, but doesn't say anything about saturation.

[jinz2014]

Do you have any update about this ?

[LU-JOHN]

The recommended workaround is to use clamp and then convert. Something like:

input = sycl::vec<float, 4>;

auto clamped = clamp(input,-128.0,127.0); auto result = clamped.template convert<char>();