[SYCL][FPGA] Add new argument to __spirv_ArbitraryFloatPowNINTEL (#5170)

Signed-off-by: Mikhail Lychkov <[email protected]>

Author: mlychkov

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -697,13 +697,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;
 

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>