kokkos icon indicating copy to clipboard operation
kokkos copied to clipboard

Round to integer math functions

Open JBludau opened this issue 1 month ago • 5 comments

fixes #8702 For FP32/FP64:

  • SYCL rint is supported, but the others are missing.
  • HIP CUDA full support

For FP16:

  • SYCL rint and round for half and bhalf.
  • HIP CUDA only rint for half and bhalf

If my derivations are correct, we could do the following for half: Implement the l ll versions of rint and round for half precision by just casting the result of to long or long long since the range of half fits inside int. For the l ll versions of rint and round for bhalf we fall back to the float ones.

JBludau avatar Dec 02 '25 21:12 JBludau

Note that most of these functions are provided in the Intel Math Function Device Libraray (https://www.intel.com/content/www/us/en/docs/dpcpp-cpp-compiler/developer-guide-reference/2025-2/imf-transcendental-math-functions.html), e.g., as sycl::ext::intel::math::lround.

masterleinad avatar Dec 02 '25 22:12 masterleinad

Note that most of these functions are provided in the Intel Math Function Device Libraray (https://www.intel.com/content/www/us/en/docs/dpcpp-cpp-compiler/developer-guide-reference/2025-2/imf-transcendental-math-functions.html), e.g., as sycl::ext::intel::math::lround.

I would propose we do the math functions for the others first and do a followup where we try to forward to whatever we can find in sycl::ext::intel::math

JBludau avatar Dec 02 '25 22:12 JBludau

What about the math functions for half precision and quarter precision types?

masterleinad avatar Dec 02 '25 23:12 masterleinad

What about the math functions for half precision and quarter precision types?

Sycl seems to support rint and even round for half and bhalf. HIP and CUDA do only supply rint for half and bhalf but not round.

If my derivations are correct, we could do the following for half:

  • Implement the l ll versions of rint and round for half precision by just casting the result of to long or long longsince the range of half fits inside int.
  • For the l ll versions of rint and round for bhalf we fall back to the float ones.

JBludau avatar Dec 03 '25 22:12 JBludau

retest this please

JBludau avatar Dec 08 '25 20:12 JBludau

SYCL CI is failing with

/home/runner/_work/kokkos/kokkos/kokkos/core/src/impl/Kokkos_Half_MathematicalFunctions.hpp:215:26: error: no member named 'impl_lround' in namespace 'Kokkos::Impl'
  215 |     return Kokkos::Impl::impl_lround(x);
      |            ~~~~~~~~~~~~~~^
/home/runner/_work/kokkos/kokkos/kokkos/core/src/impl/Kokkos_Half_MathematicalFunctions.hpp:218:26: error: no member named 'impl_llround' in namespace 'Kokkos::Impl'
  218 |     return Kokkos::Impl::impl_llround(x);
      |            ~~~~~~~~~~~~~~^
/home/runner/_work/kokkos/kokkos/kokkos/core/src/impl/Kokkos_Half_MathematicalFunctions.hpp:221:26: error: no member named 'impl_lrint' in namespace 'Kokkos::Impl'
  221 |     return Kokkos::Impl::impl_lrint(x);
      |            ~~~~~~~~~~~~~~^
Build succeeded.
/home/runner/_work/kokkos/kokkos/kokkos/core/src/impl/Kokkos_Half_MathematicalFunctions.hpp:224:26: error: no member named 'impl_llrint' in namespace 'Kokkos::Impl'
  224 |     return Kokkos::Impl::impl_llrint(x);
      |            ~~~~~~~~~~~~~~^

masterleinad avatar Dec 16 '25 20:12 masterleinad