From b72cdc99427969c0294d11446fa08113b814f510 Mon Sep 17 00:00:00 2001
From: Mikhail Lychkov <mikhail.lychkov@intel.com>
Date: Fri, 17 Dec 2021 12:05:36 +0300
Subject: [PATCH] [SYCL][FPGA] Add new argument to
 __spirv_ArbitraryFloatPowNINTEL

Signed-off-by: Mikhail Lychkov <mikhail.lychkov@intel.com>
---
 sycl/include/CL/__spirv/spirv_ops.hpp          | 8 ++++----
 sycl/test/check_device_code/fpga_ihs_float.cpp | 5 +++--
 2 files changed, 7 insertions(+), 6 deletions(-)

diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp
index 55b0764aefe1f..15fdaf03c94df 100644
--- a/sycl/include/CL/__spirv/spirv_ops.hpp
+++ b/sycl/include/CL/__spirv/spirv_ops.hpp
@@ -675,13 +675,13 @@ __spirv_ArbitraryFloatPowRINTEL(cl::sycl::detail::ap_int<WA> A, int32_t MA,
                                 int32_t RoundingAccuracy = 0) noexcept;
 
 // PowN built-in calculates `A^B` where `A` is arbitrary precision floating
-// point number and `B` is arbitrary precision integer, i.e. its width doesn't
-// depend on sum of exponent and mantissa.
+// point number and `B` is signed or unsigned arbitrary precision integer,
+// i.e. its width doesn't depend on sum of exponent and mantissa.
 template <int WA, int WB, int Wout>
 extern SYCL_EXTERNAL cl::sycl::detail::ap_int<Wout>
 __spirv_ArbitraryFloatPowNINTEL(cl::sycl::detail::ap_int<WA> A, int32_t MA,
-                                cl::sycl::detail::ap_int<WB> B, int32_t Mout,
-                                int32_t EnableSubnormals = 0,
+                                cl::sycl::detail::ap_int<WB> B, bool SignOfB,
+                                int32_t Mout, int32_t EnableSubnormals = 0,
                                 int32_t RoundingMode = 0,
                                 int32_t RoundingAccuracy = 0) noexcept;
 
diff --git a/sycl/test/check_device_code/fpga_ihs_float.cpp b/sycl/test/check_device_code/fpga_ihs_float.cpp
index ec55bd3ff0f6c..1699c7b024d67 100644
--- a/sycl/test/check_device_code/fpga_ihs_float.cpp
+++ b/sycl/test/check_device_code/fpga_ihs_float.cpp
@@ -16,6 +16,7 @@ constexpr int32_t RndMode = 2;
 constexpr int32_t RndAcc = 1;
 constexpr bool FromSign = false;
 constexpr bool ToSign = true;
+constexpr bool SignOfB = false;
 
 template <int EA, int MA, int Eout, int Mout>
 void ap_float_cast() {
@@ -390,8 +391,8 @@ void ap_float_pown() {
   sycl::detail::ap_int<WB> B;
   sycl::detail::ap_int<1 + Eout + Mout> pown_res =
       __spirv_ArbitraryFloatPowNINTEL<1 + EA + MA, WB, 1 + Eout + Mout>(
-          A, MA, B, Mout, Subnorm, RndMode, RndAcc);
-  // CHECK: call spir_func signext i15 @_Z{{[0-9]+}}__spirv_ArbitraryFloatPowNINTEL{{.*}}(i12 signext {{[%a-z0-9.]+}}, i32 7, i10 signext {{[%a-z0-9.]+}}, i32 9, i32 0, i32 2, i32 1)
+          A, MA, B, SignOfB, Mout, Subnorm, RndMode, RndAcc);
+  // CHECK: call spir_func signext i15 @_Z{{[0-9]+}}__spirv_ArbitraryFloatPowNINTEL{{.*}}(i12 signext {{[%a-z0-9.]+}}, i32 7, i10 signext {{[%a-z0-9.]+}}, i1 zeroext false, i32 9, i32 0, i32 2, i32 1)
 }
 
 template <typename name, typename Func>