diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index 4bfad035404d3..06d8ac40e09ca 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -659,7 +659,8 @@ class ESIMDIntrinDescTable { {"__spirv_ConvertBF16ToFINTEL", {a(0)}}}, {"addc", {"addc", {l(0)}}}, {"subb", {"subb", {l(0)}}}, - {"bfn", {"bfn", {a(0), a(1), a(2), t(0)}}}}; + {"bfn", {"bfn", {a(0), a(1), a(2), t(0)}}}, + {"srnd", {"srnd", {a(0), a(1)}}}}; } // clang-format on diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp index a9275bdde2708..7be20c432ebfc 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp @@ -146,6 +146,12 @@ __ESIMD_INTRIN __ESIMD_raw_vec_t(T, N) __esimd_bfn(__ESIMD_raw_vec_t(T, N) src0, __ESIMD_raw_vec_t(T, N) src1, __ESIMD_raw_vec_t(T, N) src2) __ESIMD_INTRIN_END; +template +__ESIMD_INTRIN __ESIMD_raw_vec_t(sycl::half, N) + __esimd_srnd(__ESIMD_DNS::vector_type_t src1, + __ESIMD_DNS::vector_type_t src2) + __ESIMD_INTRIN_END; + #undef __ESIMD_raw_vec_t #undef __ESIMD_cpp_vec_t diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp index 9e78bfe1d4875..e69f5686e2aa6 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp @@ -1665,6 +1665,19 @@ __ESIMD_NS::simd dp4(__ESIMD_NS::simd v1, return retv; } +/// srnd - perform stochastic rounding. +/// Supported conversions: +/// float -> half +/// Available on PVC_XT+ +/// \param src0 the operand to be rounded +/// \param src1 random number used for rounding +/// \return the converted value +template +ESIMD_INLINE __ESIMD_NS::simd +srnd(__ESIMD_NS::simd src0, __ESIMD_NS::simd src1) { + return __esimd_srnd(src0.data(), src1.data()); +} + /// @} sycl_esimd_math /// @addtogroup sycl_esimd_logical diff --git a/sycl/test-e2e/ESIMD/srnd.cpp b/sycl/test-e2e/ESIMD/srnd.cpp new file mode 100644 index 0000000000000..fe8730f4b661b --- /dev/null +++ b/sycl/test-e2e/ESIMD/srnd.cpp @@ -0,0 +1,76 @@ +//==---------------- srnd.cpp - DPC++ ESIMD srnd function test -----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM +// Exceptions. See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu-intel-pvc +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#include "esimd_test_utils.hpp" + +#include +#include + +#include + +using namespace sycl; +using namespace sycl::ext::intel::esimd; +using namespace sycl::ext::intel::experimental::esimd; + +bool test(queue &Q) { + + constexpr int N = 8; + float *Input = malloc_shared(N, Q); + half *Output = malloc_shared(N * 2, Q); + + for (int i = 0; i < N; ++i) { + float value = esimd_test::getRandomValue(); + while (fabs(value) > std::numeric_limits::max() || + fabs(value) < std::numeric_limits::min() || std::isnan(value)) + value = esimd_test::getRandomValue(); + + Input[i] = value; + } + + Q.single_task([=]() SYCL_ESIMD_KERNEL { + simd InputVector; + InputVector.copy_from(Input); + { + simd RandomVector(0); + simd OutputVector = srnd(InputVector, RandomVector); + OutputVector.copy_to(Output); + } + { + simd RandomVector(0xFFFF); + simd OutputVector = srnd(InputVector, RandomVector); + OutputVector.copy_to(Output + N); + } + }).wait(); + bool ReturnValue = true; + for (int i = 0; i < N; ++i) { + if (Output[i + N] == Output[i]) { + ReturnValue = false; + break; + } + } + + free(Input, Q); + free(Output, Q); + return ReturnValue; +} + +// --- The entry point. + +int main(void) { + queue Q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler()); + esimd_test::printTestLabel(Q); + bool Pass = true; + + Pass &= test(Q); + + std::cout << (Pass ? "Test Passed\n" : "Test FAILED\n"); + return Pass ? 0 : 1; +}