[SYCL][ESIMD] Add frem function (#15117)

Add support for the frem function.

Support for this was just added to IGC, so we aren't using the driver in
CI yet, so the test won't actually run.

I manually tested single type which works with the version in
`REQUIRES-INTEL-DRIVER`, but I found an issue in the driver with double
so it's disabled in the test for now.

Signed-off-by: Sarnie, Nick <[email protected]>

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp

@@ -138,6 +138,12 @@ template <typename T, int N>
 extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
     __spirv_ocl_clz(__ESIMD_raw_vec_t(T, N) src0) __ESIMD_INTRIN_END;
 
+template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_FRem(T);
+template <typename T, int N>
+extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
+    __spirv_FRem(__ESIMD_raw_vec_t(T, N) src0,
+                 __ESIMD_raw_vec_t(T, N) src1) __ESIMD_INTRIN_END;
+
 #undef __ESIMD_raw_vec_t
 #undef __ESIMD_cpp_vec_t
 

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

@@ -721,6 +721,18 @@ srnd(__ESIMD_NS::simd<float, N> src0, __ESIMD_NS::simd<uint16_t, N> src1) {
   return __esimd_srnd<N>(src0.data(), src1.data());
 }
 
+/// frem - compute the remainder from floating point division.
+/// \param src0 the first operand to be used for division.
+/// \param src1 the second operand to be used for division.
+/// \return the remainder from the division.
+template <typename T, int N>
+ESIMD_INLINE __ESIMD_NS::simd<T, N> frem(__ESIMD_NS::simd<T, N> src0,
+                                         __ESIMD_NS::simd<T, N> src1) {
+  static_assert(std::is_same_v<T, float> || std::is_same_v<T, double>,
+                "Element type must be float or double");
+  return __spirv_FRem<T, N>(src0.data(), src1.data());
+}
+
 /// @} sycl_esimd_math
 
 /// @addtogroup sycl_esimd_logical

sycl/test-e2e/ESIMD/frem.cpp

@@ -0,0 +1,101 @@
+//==---------------- frem.cpp - DPC++ ESIMD on-device 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-INTEL-DRIVER: lin: 30623
+// RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %{run} %t.out
+
+#include "esimd_test_utils.hpp"
+
+#include <cmath>
+#include <sycl/ext/intel/experimental/esimd/math.hpp>
+
+using namespace sycl;
+
+template <typename T> bool test(queue q) {
+  std::cout << "Running case: T=" << esimd_test::type_name<T>() << std::endl;
+  constexpr unsigned Size = 16;
+  constexpr unsigned VL = 16;
+
+  T *A = new T[Size];
+  T *B = new T[Size];
+  T *C = new T[Size];
+
+  for (unsigned i = 0; i < Size; ++i) {
+    A[i] = i + 500;
+    B[i] = i + 1;
+    C[i] = 0.0f;
+  }
+
+  try {
+    buffer<T, 1> bufa(A, range<1>(Size));
+    buffer<T, 1> bufb(B, range<1>(Size));
+    buffer<T, 1> bufc(C, range<1>(Size));
+
+    auto e = q.submit([&](handler &cgh) {
+      auto PA = bufa.template get_access<access::mode::read>(cgh);
+      auto PB = bufb.template get_access<access::mode::read>(cgh);
+      auto PC = bufc.template get_access<access::mode::write>(cgh);
+      cgh.single_task([=]() SYCL_ESIMD_KERNEL {
+        using namespace sycl::ext::intel::esimd;
+        unsigned int offset = 0;
+        simd<T, VL> va;
+        va.copy_from(PA, offset);
+        simd<T, VL> vb;
+        vb.copy_from(PB, offset);
+        simd<T, VL> vc = ext::intel::experimental::esimd::frem(va, vb);
+        vc.copy_to(PC, offset);
+      });
+    });
+    e.wait();
+  } catch (sycl::exception const &e) {
+    std::cout << "SYCL exception caught: " << e.what() << '\n';
+
+    delete[] A;
+    delete[] B;
+    delete[] C;
+    return 0;
+  }
+
+  int err_cnt = 0;
+
+  for (unsigned i = 0; i < Size; ++i) {
+    T expected = std::remainder(A[i], B[i]);
+    if (C[i] != expected) {
+      if (++err_cnt < 10) {
+        std::cout << "failed at index " << i << ", " << std::to_string(C[i])
+                  << " != " << std::to_string(expected) << "\n";
+      }
+    }
+  }
+  if (err_cnt > 0) {
+    std::cout << "  pass rate: "
+              << ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% ("
+              << (Size - err_cnt) << "/" << Size << ")\n";
+  }
+
+  delete[] A;
+  delete[] B;
+  delete[] C;
+
+  std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n");
+  return err_cnt == 0;
+}
+
+int main() {
+  auto q = queue{gpu_selector_v};
+  esimd_test::printTestLabel(q);
+  bool passed = true;
+
+  passed &= test<float>(q);
+  // TODO: Enable when driver issue fixed
+#if 0
+  if (q.get_device().has(sycl::aspect::fp64))
+    passed &= test<double>(q);
+#endif
+  return passed ? 0 : 1;
+}