[SYCL][ESIMD] Add srnd API to ESIMD (#12167)

Author: fineg74

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
 

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 <int N>
+__ESIMD_INTRIN __ESIMD_raw_vec_t(sycl::half, N)
+    __esimd_srnd(__ESIMD_DNS::vector_type_t<float, N> src1,
+                 __ESIMD_DNS::vector_type_t<uint16_t, N> src2)
+        __ESIMD_INTRIN_END;
+
 #undef __ESIMD_raw_vec_t
 #undef __ESIMD_cpp_vec_t
 

sycl/include/sycl/ext/intel/experimental/esimd/math.hpp

@@ -1665,6 +1665,19 @@ __ESIMD_NS::simd<T, N> dp4(__ESIMD_NS::simd<T, N> 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 <int N>
+ESIMD_INLINE __ESIMD_NS::simd<sycl::half, N>
+srnd(__ESIMD_NS::simd<float, N> src0, __ESIMD_NS::simd<uint16_t, N> src1) {
+  return __esimd_srnd<N>(src0.data(), src1.data());
+}
+
 /// @} sycl_esimd_math
 
 /// @addtogroup sycl_esimd_logical

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 <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+#include <iostream>
+
+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<float>(N, Q);
+  half *Output = malloc_shared<half>(N * 2, Q);
+
+  for (int i = 0; i < N; ++i) {
+    float value = esimd_test::getRandomValue<float>();
+    while (fabs(value) > std::numeric_limits<half>::max() ||
+           fabs(value) < std::numeric_limits<half>::min() || std::isnan(value))
+      value = esimd_test::getRandomValue<float>();
+
+    Input[i] = value;
+  }
+
+  Q.single_task([=]() SYCL_ESIMD_KERNEL {
+     simd<float, N> InputVector;
+     InputVector.copy_from(Input);
+     {
+       simd<uint16_t, N> RandomVector(0);
+       simd<half, N> OutputVector = srnd(InputVector, RandomVector);
+       OutputVector.copy_to(Output);
+     }
+     {
+       simd<uint16_t, N> RandomVector(0xFFFF);
+       simd<half, N> 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;
+}