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 bbf565e0654a4..c156f7381cabb 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 @@ -138,6 +138,12 @@ template extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N) __spirv_ocl_clz(__ESIMD_raw_vec_t(T, N) src0) __ESIMD_INTRIN_END; +template extern __DPCPP_SYCL_EXTERNAL T __spirv_FRem(T); +template +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 diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp index 99674732680a4..eb1c825a6d03d 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp @@ -721,6 +721,18 @@ srnd(__ESIMD_NS::simd src0, __ESIMD_NS::simd src1) { return __esimd_srnd(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 +ESIMD_INLINE __ESIMD_NS::simd frem(__ESIMD_NS::simd src0, + __ESIMD_NS::simd src1) { + static_assert(std::is_same_v || std::is_same_v, + "Element type must be float or double"); + return __spirv_FRem(src0.data(), src1.data()); +} + /// @} sycl_esimd_math /// @addtogroup sycl_esimd_logical diff --git a/sycl/test-e2e/ESIMD/frem.cpp b/sycl/test-e2e/ESIMD/frem.cpp new file mode 100644 index 0000000000000..7f69ac6730e54 --- /dev/null +++ b/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 +#include + +using namespace sycl; + +template bool test(queue q) { + std::cout << "Running case: T=" << esimd_test::type_name() << 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 bufa(A, range<1>(Size)); + buffer bufb(B, range<1>(Size)); + buffer bufc(C, range<1>(Size)); + + auto e = q.submit([&](handler &cgh) { + auto PA = bufa.template get_access(cgh); + auto PB = bufb.template get_access(cgh); + auto PC = bufc.template get_access(cgh); + cgh.single_task([=]() SYCL_ESIMD_KERNEL { + using namespace sycl::ext::intel::esimd; + unsigned int offset = 0; + simd va; + va.copy_from(PA, offset); + simd vb; + vb.copy_from(PB, offset); + simd 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(q); + // TODO: Enable when driver issue fixed +#if 0 + if (q.get_device().has(sycl::aspect::fp64)) + passed &= test(q); +#endif + return passed ? 0 : 1; +}