diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index c8579d3f49a2e..55b0764aefe1f 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -325,281 +325,363 @@ __SYCL_CONVERGENT__ extern SYCL_EXTERNAL void __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint64_t *Ptr, dataT Data) noexcept; template -extern SYCL_EXTERNAL ap_int -__spirv_FixedSqrtINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, int32_t Overflow = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_FixedSqrtINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, + int32_t Overflow = 0) noexcept; template -extern SYCL_EXTERNAL ap_int -__spirv_FixedRecipINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_FixedRecipINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, int32_t Overflow = 0) noexcept; template -extern SYCL_EXTERNAL ap_int -__spirv_FixedRsqrtINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_FixedRsqrtINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, int32_t Overflow = 0) noexcept; template -extern SYCL_EXTERNAL ap_int -__spirv_FixedSinINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, int32_t Overflow = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_FixedSinINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, + int32_t Overflow = 0) noexcept; template -extern SYCL_EXTERNAL ap_int -__spirv_FixedCosINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, int32_t Overflow = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_FixedCosINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, + int32_t Overflow = 0) noexcept; template -extern SYCL_EXTERNAL ap_int<2 * rW> -__spirv_FixedSinCosINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, +extern SYCL_EXTERNAL cl::sycl::detail::ap_int<2 * rW> +__spirv_FixedSinCosINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, int32_t Overflow = 0) noexcept; template -extern SYCL_EXTERNAL ap_int -__spirv_FixedSinPiINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_FixedSinPiINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, int32_t Overflow = 0) noexcept; template -extern SYCL_EXTERNAL ap_int -__spirv_FixedCosPiINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_FixedCosPiINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, int32_t Overflow = 0) noexcept; template -extern SYCL_EXTERNAL ap_int<2 * rW> -__spirv_FixedSinCosPiINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, +extern SYCL_EXTERNAL cl::sycl::detail::ap_int<2 * rW> +__spirv_FixedSinCosPiINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, int32_t Overflow = 0) noexcept; template -extern SYCL_EXTERNAL ap_int -__spirv_FixedLogINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, int32_t Overflow = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_FixedLogINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, + int32_t Overflow = 0) noexcept; template -extern SYCL_EXTERNAL ap_int -__spirv_FixedExpINTEL(ap_int a, bool S, int32_t I, int32_t rI, - int32_t Quantization = 0, int32_t Overflow = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_FixedExpINTEL(cl::sycl::detail::ap_int a, bool S, int32_t I, + int32_t rI, int32_t Quantization = 0, + int32_t Overflow = 0) noexcept; // In the following built-ins width of arbitrary precision integer type for // a floating point variable should be equal to sum of corresponding // exponent width E, mantissa width M and 1 for sign bit. I.e. WA = EA + MA + 1. template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatCastINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatCastINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatCastFromIntINTEL( - ap_int A, int32_t Mout, bool FromSign = false, - int32_t EnableSubnormals = 0, int32_t RoundingMode = 0, - int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatCastFromIntINTEL(cl::sycl::detail::ap_int A, + int32_t Mout, bool FromSign = false, + int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatCastToIntINTEL( - ap_int A, int32_t MA, bool ToSign = false, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatCastToIntINTEL(cl::sycl::detail::ap_int A, int32_t MA, + bool ToSign = false, + int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatAddINTEL( - ap_int A, int32_t MA, ap_int B, int32_t MB, int32_t Mout, - int32_t EnableSubnormals = 0, int32_t RoundingMode = 0, - int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatAddINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatSubINTEL( - ap_int A, int32_t MA, ap_int B, int32_t MB, int32_t Mout, - int32_t EnableSubnormals = 0, int32_t RoundingMode = 0, - int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatSubINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatMulINTEL( - ap_int A, int32_t MA, ap_int B, int32_t MB, int32_t Mout, - int32_t EnableSubnormals = 0, int32_t RoundingMode = 0, - int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatMulINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatDivINTEL( - ap_int A, int32_t MA, ap_int B, int32_t MB, int32_t Mout, - int32_t EnableSubnormals = 0, int32_t RoundingMode = 0, - int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatDivINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; // Comparison built-ins don't use Subnormal Support, Rounding Mode and // Rounding Accuracy. template extern SYCL_EXTERNAL bool -__spirv_ArbitraryFloatGTINTEL(ap_int A, int32_t MA, ap_int B, +__spirv_ArbitraryFloatGTINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB) noexcept; template extern SYCL_EXTERNAL bool -__spirv_ArbitraryFloatGEINTEL(ap_int A, int32_t MA, ap_int B, +__spirv_ArbitraryFloatGEINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB) noexcept; template extern SYCL_EXTERNAL bool -__spirv_ArbitraryFloatLTINTEL(ap_int A, int32_t MA, ap_int B, +__spirv_ArbitraryFloatLTINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB) noexcept; template extern SYCL_EXTERNAL bool -__spirv_ArbitraryFloatLEINTEL(ap_int A, int32_t MA, ap_int B, +__spirv_ArbitraryFloatLEINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB) noexcept; template extern SYCL_EXTERNAL bool -__spirv_ArbitraryFloatEQINTEL(ap_int A, int32_t MA, ap_int B, +__spirv_ArbitraryFloatEQINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatRecipINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatRecipINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatRSqrtINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatRSqrtINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatCbrtINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatCbrtINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatHypotINTEL( - ap_int A, int32_t MA, ap_int B, int32_t MB, int32_t Mout, - int32_t EnableSubnormals = 0, int32_t RoundingMode = 0, - int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatHypotINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatSqrtINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatSqrtINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatLogINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatLogINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatLog2INTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatLog2INTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatLog10INTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatLog10INTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatLog1pINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatLog1pINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatExpINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatExpINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatExp2INTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatExp2INTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatExp10INTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatExp10INTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatExpm1INTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatExpm1INTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatSinINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatSinINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatCosINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatCosINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; // Result value contains both values of sine and cosine and so has the size of // 2 * Wout where Wout is equal to (1 + Eout + Mout). template -extern SYCL_EXTERNAL ap_int<2 * Wout> __spirv_ArbitraryFloatSinCosINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int<2 * Wout> +__spirv_ArbitraryFloatSinCosINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatSinPiINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatSinPiINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatCosPiINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatCosPiINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; // Result value contains both values of sine(A*pi) and cosine(A*pi) and so has // the size of 2 * Wout where Wout is equal to (1 + Eout + Mout). template -extern SYCL_EXTERNAL ap_int<2 * Wout> __spirv_ArbitraryFloatSinCosPiINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int<2 * Wout> +__spirv_ArbitraryFloatSinCosPiINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatASinINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatASinINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatASinPiINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatASinPiINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatACosINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatACosINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatACosPiINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatACosPiINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatATanINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatATanINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatATanPiINTEL( - ap_int A, int32_t MA, int32_t Mout, int32_t EnableSubnormals = 0, - int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatATanPiINTEL(cl::sycl::detail::ap_int A, int32_t MA, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatATan2INTEL( - ap_int A, int32_t MA, ap_int B, int32_t MB, int32_t Mout, - int32_t EnableSubnormals = 0, int32_t RoundingMode = 0, - int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatATan2INTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatPowINTEL( - ap_int A, int32_t MA, ap_int B, int32_t MB, int32_t Mout, - int32_t EnableSubnormals = 0, int32_t RoundingMode = 0, - int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatPowINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + int32_t RoundingAccuracy = 0) noexcept; template -extern SYCL_EXTERNAL ap_int __spirv_ArbitraryFloatPowRINTEL( - ap_int A, int32_t MA, ap_int B, int32_t MB, int32_t Mout, - int32_t EnableSubnormals = 0, int32_t RoundingMode = 0, - int32_t RoundingAccuracy = 0) noexcept; +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatPowRINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t MB, + int32_t Mout, int32_t EnableSubnormals = 0, + int32_t RoundingMode = 0, + 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. template -extern SYCL_EXTERNAL ap_int -__spirv_ArbitraryFloatPowNINTEL(ap_int A, int32_t MA, ap_int B, - int32_t Mout, int32_t EnableSubnormals = 0, +extern SYCL_EXTERNAL cl::sycl::detail::ap_int +__spirv_ArbitraryFloatPowNINTEL(cl::sycl::detail::ap_int A, int32_t MA, + cl::sycl::detail::ap_int B, int32_t Mout, + int32_t EnableSubnormals = 0, int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; diff --git a/sycl/include/CL/__spirv/spirv_types.hpp b/sycl/include/CL/__spirv/spirv_types.hpp index 63c2f26ca7387..cae0f38c11748 100644 --- a/sycl/include/CL/__spirv/spirv_types.hpp +++ b/sycl/include/CL/__spirv/spirv_types.hpp @@ -8,6 +8,8 @@ #pragma once +#include + #include #include @@ -133,8 +135,14 @@ struct ConstantPipeStorage { int32_t _Capacity; }; +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { // Arbitrary precision integer type template using ap_int = _ExtInt(Bits); +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) #endif // __SYCL_DEVICE_ONLY__ // This class does not have definition, it is only predeclared here. diff --git a/sycl/test/check_device_code/ap_fixed.cpp b/sycl/test/check_device_code/ap_fixed.cpp index 00747784946e9..2833ded76a35a 100644 --- a/sycl/test/check_device_code/ap_fixed.cpp +++ b/sycl/test/check_device_code/ap_fixed.cpp @@ -12,77 +12,77 @@ template void sqrt() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_Sqrt = __spirv_FixedSqrtINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func signext i5 @_Z[[#]]__spirv_FixedSqrtINTEL{{.*}}(i13 signext %[[#]], i1 zeroext false, i32 2, i32 2, i32 0, i32 0) } template void recip() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_Recip = __spirv_FixedRecipINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func signext i8 @_Z[[#]]__spirv_FixedRecipINTEL{{.*}}(i3 signext %[[#]], i1 zeroext true, i32 4, i32 4, i32 0, i32 0) } template void rsqrt() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_Rsqrt = __spirv_FixedRsqrtINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func signext i10 @_Z[[#]]__spirv_FixedRsqrtINTEL{{.*}}(i11 signext %[[#]], i1 zeroext false, i32 8, i32 6, i32 0, i32 0) } template void sin() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_Sin = __spirv_FixedSinINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func signext i11 @_Z[[#]]__spirv_FixedSinINTEL{{.*}}(i17 signext %[[#]], i1 zeroext true, i32 7, i32 5, i32 0, i32 0) } template void cos() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_Cos = __spirv_FixedCosINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func signext i28 @_Z[[#]]__spirv_FixedCosINTEL{{.*}}(i35 %[[#]], i1 zeroext false, i32 9, i32 3, i32 0, i32 0) } template void sin_cos() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_SinCos = __spirv_FixedSinCosINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func i40 @_Z[[#]]__spirv_FixedSinCosINTEL{{.*}}(i31 signext %[[#]], i1 zeroext true, i32 10, i32 12, i32 0, i32 0) } template void sin_pi() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_SinPi = __spirv_FixedSinPiINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func signext i5 @_Z[[#]]__spirv_FixedSinPiINTEL{{.*}}(i60 %[[#]], i1 zeroext false, i32 2, i32 2, i32 0, i32 0) } template void cos_pi() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_CosPi = __spirv_FixedCosPiINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func signext i16 @_Z[[#]]__spirv_FixedCosPiINTEL{{.*}}(i28 signext %[[#]], i1 zeroext false, i32 8, i32 5, i32 0, i32 0) } template void sin_cos_pi() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_SinCosPi = __spirv_FixedSinCosPiINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func signext i10 @_Z[[#]]__spirv_FixedSinCosPiINTEL{{.*}}(i13 signext %[[#]], i1 zeroext false, i32 2, i32 2, i32 0, i32 0) } template void log() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_Log = __spirv_FixedLogINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func i44 @_Z[[#]]__spirv_FixedLogINTEL{{.*}}(i64 %[[#]], i1 zeroext true, i32 24, i32 22, i32 0, i32 0) } template void exp() { - ap_int a; + sycl::detail::ap_int a; auto ap_fixed_Exp = __spirv_FixedExpINTEL(a, S, I, rI); // CHECK: %{{.*}} = call spir_func i34 @_Z[[#]]__spirv_FixedExpINTEL{{.*}}(i44 %[[#]], i1 zeroext false, i32 20, i32 20, i32 0, i32 0) } diff --git a/sycl/test/check_device_code/fpga_ihs_float.cpp b/sycl/test/check_device_code/fpga_ihs_float.cpp index 0b16ab8fe508c..ec55bd3ff0f6c 100644 --- a/sycl/test/check_device_code/fpga_ihs_float.cpp +++ b/sycl/test/check_device_code/fpga_ihs_float.cpp @@ -19,8 +19,8 @@ constexpr bool ToSign = true; template void ap_float_cast() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> float_cast_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> float_cast_res = __spirv_ArbitraryFloatCastINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i40 @_Z{{[0-9]+}}__spirv_ArbitraryFloatCastINTEL{{.*}}(i40 {{[%a-z0-9.]+}}, i32 28, i32 30, i32 0, i32 2, i32 1) @@ -28,8 +28,8 @@ void ap_float_cast() { template void ap_float_cast_from_int() { - ap_int A; - ap_int<1 + Eout + Mout> cast_from_int_res = + sycl::detail::ap_int A; + sycl::detail::ap_int<1 + Eout + Mout> cast_from_int_res = __spirv_ArbitraryFloatCastFromIntINTEL( A, Mout, FromSign, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i25 @_Z{{[0-9]+}}__spirv_ArbitraryFloatCastFromIntINTEL{{.*}}(i43 {{[%a-z0-9.]+}}, i32 16, i1 zeroext false, i32 0, i32 2, i32 1) @@ -37,8 +37,8 @@ void ap_float_cast_from_int() { template void ap_float_cast_to_int() { - ap_int<1 + EA + MA> A; - ap_int cast_to_int_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int cast_to_int_res = __spirv_ArbitraryFloatCastToIntINTEL<1 + EA + MA, Wout>( A, MA, ToSign, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i30 @_Z{{[0-9]+}}__spirv_ArbitraryFloatCastToIntINTEL{{.*}}(i23 signext {{[%a-z0-9.]+}}, i32 15, i1 zeroext true, i32 0, i32 2, i32 1) @@ -46,9 +46,9 @@ void ap_float_cast_to_int() { template void ap_float_add() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; - ap_int<1 + Eout + Mout> add_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + Eout + Mout> add_res = __spirv_ArbitraryFloatAddINTEL<1 + EA + MA, 1 + EB + MB, 1 + Eout + Mout>( A, MA, B, MB, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i14 @_Z{{[0-9]+}}__spirv_ArbitraryFloatAddINTEL{{.*}}(i13 signext {{[%a-z0-9.]+}}, i32 7, i15 signext {{[%a-z0-9.]+}}, i32 8, i32 9, i32 0, i32 2, i32 1) @@ -57,9 +57,9 @@ void ap_float_add() { template void ap_float_sub() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; - ap_int<1 + Eout + Mout> sub_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + Eout + Mout> sub_res = __spirv_ArbitraryFloatSubINTEL<1 + EA + MA, 1 + EB + MB, 1 + Eout + Mout>( A, MA, B, MB, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i13 @_Z{{[0-9]+}}__spirv_ArbitraryFloatSubINTEL{{.*}}(i9 signext {{[%a-z0-9.]+}}, i32 4, i11 signext {{[%a-z0-9.]+}}, i32 5, i32 6, i32 0, i32 2, i32 1) @@ -67,9 +67,9 @@ void ap_float_sub() { template void ap_float_mul() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; - ap_int<1 + Eout + Mout> mul_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + Eout + Mout> mul_res = __spirv_ArbitraryFloatMulINTEL<1 + EA + MA, 1 + EB + MB, 1 + Eout + Mout>( A, MA, B, MB, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i51 @_Z{{[0-9]+}}__spirv_ArbitraryFloatMulINTEL{{.*}}(i51 {{[%a-z0-9.]+}}, i32 34, i51 {{[%a-z0-9.]+}}, i32 34, i32 34, i32 0, i32 2, i32 1) @@ -77,9 +77,9 @@ void ap_float_mul() { template void ap_float_div() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; - ap_int<1 + Eout + Mout> div_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + Eout + Mout> div_res = __spirv_ArbitraryFloatDivINTEL<1 + EA + MA, 1 + EB + MB, 1 + Eout + Mout>( A, MA, B, MB, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i18 @_Z{{[0-9]+}}__spirv_ArbitraryFloatDivINTEL{{.*}}(i16 signext {{[%a-z0-9.]+}}, i32 11, i16 signext {{[%a-z0-9.]+}}, i32 11, i32 12, i32 0, i32 2, i32 1) @@ -87,48 +87,48 @@ void ap_float_div() { template void ap_float_gt() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; bool gt_res = __spirv_ArbitraryFloatGTINTEL<1 + EA + MA, 1 + EB + MB>(A, MA, B, MB); // CHECK: call spir_func zeroext i1 @_Z{{[0-9]+}}__spirv_ArbitraryFloatGTINTEL{{.*}}(i63 {{[%a-z0-9.]+}}, i32 42, i63 {{[%a-z0-9.]+}}, i32 41) } template void ap_float_ge() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; bool ge_res = __spirv_ArbitraryFloatGEINTEL<1 + EA + MA, 1 + EB + MB>(A, MA, B, MB); // CHECK: call spir_func zeroext i1 @_Z{{[0-9]+}}__spirv_ArbitraryFloatGEINTEL{{.*}}(i47 {{[%a-z0-9.]+}}, i32 27, i47 {{[%a-z0-9.]+}}, i32 27) } template void ap_float_lt() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; bool lt_res = __spirv_ArbitraryFloatLTINTEL<1 + EA + MA, 1 + EB + MB>(A, MA, B, MB); // CHECK: call spir_func zeroext i1 @_Z{{[0-9]+}}__spirv_ArbitraryFloatLTINTEL{{.*}}(i5 signext {{[%a-z0-9.]+}}, i32 2, i7 signext {{[%a-z0-9.]+}}, i32 3) } template void ap_float_le() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; bool le_res = __spirv_ArbitraryFloatLEINTEL<1 + EA + MA, 1 + EB + MB>(A, MA, B, MB); // CHECK: call spir_func zeroext i1 @_Z{{[0-9]+}}__spirv_ArbitraryFloatLEINTEL{{.*}}(i55 {{[%a-z0-9.]+}}, i32 27, i55 {{[%a-z0-9.]+}}, i32 28) } template void ap_float_eq() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; bool eq_res = __spirv_ArbitraryFloatEQINTEL<1 + EA + MA, 1 + EB + MB>(A, MA, B, MB); // CHECK: call spir_func zeroext i1 @_Z{{[0-9]+}}__spirv_ArbitraryFloatEQINTEL{{.*}}(i20 signext {{[%a-z0-9.]+}}, i32 12, i15 signext {{[%a-z0-9.]+}}, i32 7) } template void ap_float_recip() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> recip_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> recip_res = __spirv_ArbitraryFloatRecipINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i39 @_Z{{[0-9]+}}__spirv_ArbitraryFloatRecipINTEL{{.*}}(i39 {{[%a-z0-9.]+}}, i32 29, i32 29, i32 0, i32 2, i32 1) @@ -136,8 +136,8 @@ void ap_float_recip() { template void ap_float_rsqrt() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> rsqrt_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> rsqrt_res = __spirv_ArbitraryFloatRSqrtINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i34 @_Z{{[0-9]+}}__spirv_ArbitraryFloatRSqrtINTEL{{.*}}(i32 {{[%a-z0-9.]+}}, i32 19, i32 20, i32 0, i32 2, i32 1) @@ -145,8 +145,8 @@ void ap_float_rsqrt() { template void ap_float_cbrt() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> cbrt_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> cbrt_res = __spirv_ArbitraryFloatCbrtINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i2 @_Z{{[0-9]+}}__spirv_ArbitraryFloatCbrtINTEL{{.*}}(i2 signext {{[%a-z0-9.]+}}, i32 1, i32 1, i32 0, i32 2, i32 1) @@ -154,18 +154,19 @@ void ap_float_cbrt() { template void ap_float_hypot() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; - ap_int<1 + Eout + Mout> hypot_res = - __spirv_ArbitraryFloatHypotINTEL<1 + EA + MA, 1 + EB + MB, 1 + Eout + Mout>( + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + Eout + Mout> hypot_res = + __spirv_ArbitraryFloatHypotINTEL<1 + EA + MA, 1 + EB + MB, + 1 + Eout + Mout>( A, MA, B, MB, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i42 @_Z{{[0-9]+}}__spirv_ArbitraryFloatHypotINTEL{{.*}}(i41 {{[%a-z0-9.]+}}, i32 20, i43 {{[%a-z0-9.]+}}, i32 21, i32 22, i32 0, i32 2, i32 1) } template void ap_float_sqrt() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> sqrt_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> sqrt_res = __spirv_ArbitraryFloatSqrtINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i17 @_Z{{[0-9]+}}__spirv_ArbitraryFloatSqrtINTEL{{.*}}(i15 signext {{[%a-z0-9.]+}}, i32 7, i32 8, i32 0, i32 2, i32 1) @@ -173,8 +174,8 @@ void ap_float_sqrt() { template void ap_float_log() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> log_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> log_res = __spirv_ArbitraryFloatLogINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i50 @_Z{{[0-9]+}}__spirv_ArbitraryFloatLogINTEL{{.*}}(i50 {{[%a-z0-9.]+}}, i32 19, i32 30, i32 0, i32 2, i32 1) @@ -182,8 +183,8 @@ void ap_float_log() { template void ap_float_log2() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> log2_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> log2_res = __spirv_ArbitraryFloatLog2INTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i38 @_Z{{[0-9]+}}__spirv_ArbitraryFloatLog2INTEL{{.*}}(i38 {{[%a-z0-9.]+}}, i32 20, i32 19, i32 0, i32 2, i32 1) @@ -191,8 +192,8 @@ void ap_float_log2() { template void ap_float_log10() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> log10_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> log10_res = __spirv_ArbitraryFloatLog10INTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i10 @_Z{{[0-9]+}}__spirv_ArbitraryFloatLog10INTEL{{.*}}(i8 signext {{[%a-z0-9.]+}}, i32 3, i32 5, i32 0, i32 2, i32 1) @@ -200,8 +201,8 @@ void ap_float_log10() { template void ap_float_log1p() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> log1p_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> log1p_res = __spirv_ArbitraryFloatLog1pINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i49 @_Z{{[0-9]+}}__spirv_ArbitraryFloatLog1pINTEL{{.*}}(i48 {{[%a-z0-9.]+}}, i32 30, i32 30, i32 0, i32 2, i32 1) @@ -209,8 +210,8 @@ void ap_float_log1p() { template void ap_float_exp() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> exp_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> exp_res = __spirv_ArbitraryFloatExpINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i42 @_Z{{[0-9]+}}__spirv_ArbitraryFloatExpINTEL{{.*}}(i42 {{[%a-z0-9.]+}}, i32 25, i32 25, i32 0, i32 2, i32 1) @@ -218,8 +219,8 @@ void ap_float_exp() { template void ap_float_exp2() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> exp2_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> exp2_res = __spirv_ArbitraryFloatExp2INTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i5 @_Z{{[0-9]+}}__spirv_ArbitraryFloatExp2INTEL{{.*}}(i3 signext {{[%a-z0-9.]+}}, i32 1, i32 2, i32 0, i32 2, i32 1) @@ -227,8 +228,8 @@ void ap_float_exp2() { template void ap_float_exp10() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> exp10_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> exp10_res = __spirv_ArbitraryFloatExp10INTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i25 @_Z{{[0-9]+}}__spirv_ArbitraryFloatExp10INTEL{{.*}}(i25 signext {{[%a-z0-9.]+}}, i32 16, i32 16, i32 0, i32 2, i32 1) @@ -236,8 +237,8 @@ void ap_float_exp10() { template void ap_float_expm1() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> expm1_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> expm1_res = __spirv_ArbitraryFloatExpm1INTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i62 @_Z{{[0-9]+}}__spirv_ArbitraryFloatExpm1INTEL{{.*}}(i64 {{[%a-z0-9.]+}}, i32 42, i32 41, i32 0, i32 2, i32 1) @@ -245,8 +246,8 @@ void ap_float_expm1() { template void ap_float_sin() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> sin_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> sin_res = __spirv_ArbitraryFloatSinINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i34 @_Z{{[0-9]+}}__spirv_ArbitraryFloatSinINTEL{{.*}}(i30 signext {{[%a-z0-9.]+}}, i32 15, i32 17, i32 0, i32 2, i32 1) @@ -254,8 +255,8 @@ void ap_float_sin() { template void ap_float_cos() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> cos_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> cos_res = __spirv_ArbitraryFloatCosINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i4 @_Z{{[0-9]+}}__spirv_ArbitraryFloatCosINTEL{{.*}}(i4 signext {{[%a-z0-9.]+}}, i32 2, i32 1, i32 0, i32 2, i32 1) @@ -263,8 +264,8 @@ void ap_float_cos() { template void ap_float_sincos() { - ap_int<1 + EA + MA> A; - ap_int<2 * (1 + Eout + Mout)> sincos_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<2 * (1 + Eout + Mout)> sincos_res = __spirv_ArbitraryFloatSinCosINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i62 @_Z{{[0-9]+}}__spirv_ArbitraryFloatSinCosINTEL{{.*}}(i27 signext {{[%a-z0-9.]+}}, i32 18, i32 20, i32 0, i32 2, i32 1) @@ -272,8 +273,8 @@ void ap_float_sincos() { template void ap_float_sinpi() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> sinpi_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> sinpi_res = __spirv_ArbitraryFloatSinPiINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i13 @_Z{{[0-9]+}}__spirv_ArbitraryFloatSinPiINTEL{{.*}}(i10 signext {{[%a-z0-9.]+}}, i32 6, i32 6, i32 0, i32 2, i32 1) @@ -281,8 +282,8 @@ void ap_float_sinpi() { template void ap_float_cospi() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> cospi_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> cospi_res = __spirv_ArbitraryFloatCosPiINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i59 @_Z{{[0-9]+}}__spirv_ArbitraryFloatCosPiINTEL{{.*}}(i59 {{[%a-z0-9.]+}}, i32 40, i32 40, i32 0, i32 2, i32 1) @@ -290,8 +291,8 @@ void ap_float_cospi() { template void ap_float_sincospi() { - ap_int<1 + EA + MA> A; - ap_int<2 * (1 + Eout + Mout)> sincos_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<2 * (1 + Eout + Mout)> sincos_res = __spirv_ArbitraryFloatSinCosPiINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i64 @_Z{{[0-9]+}}__spirv_ArbitraryFloatSinCosPiINTEL{{.*}}(i30 signext {{[%a-z0-9.]+}}, i32 20, i32 20, i32 0, i32 2, i32 1) @@ -299,8 +300,8 @@ void ap_float_sincospi() { template void ap_float_asin() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> asin_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> asin_res = __spirv_ArbitraryFloatASinINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i11 @_Z{{[0-9]+}}__spirv_ArbitraryFloatASinINTEL{{.*}}(i7 signext {{[%a-z0-9.]+}}, i32 4, i32 8, i32 0, i32 2, i32 1) @@ -308,8 +309,8 @@ void ap_float_asin() { template void ap_float_asinpi() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> asinpi_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> asinpi_res = __spirv_ArbitraryFloatASinPiINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i35 @_Z{{[0-9]+}}__spirv_ArbitraryFloatASinPiINTEL{{.*}}(i35 {{[%a-z0-9.]+}}, i32 23, i32 23, i32 0, i32 2, i32 1) @@ -317,8 +318,8 @@ void ap_float_asinpi() { template void ap_float_acos() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> acos_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> acos_res = __spirv_ArbitraryFloatACosINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i14 @_Z{{[0-9]+}}__spirv_ArbitraryFloatACosINTEL{{.*}}(i14 signext {{[%a-z0-9.]+}}, i32 9, i32 10, i32 0, i32 2, i32 1) @@ -326,8 +327,8 @@ void ap_float_acos() { template void ap_float_acospi() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> acospi_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> acospi_res = __spirv_ArbitraryFloatACosPiINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i8 @_Z{{[0-9]+}}__spirv_ArbitraryFloatACosPiINTEL{{.*}}(i8 signext {{[%a-z0-9.]+}}, i32 5, i32 4, i32 0, i32 2, i32 1) @@ -335,8 +336,8 @@ void ap_float_acospi() { template void ap_float_atan() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> atan_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> atan_res = __spirv_ArbitraryFloatATanINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i44 @_Z{{[0-9]+}}__spirv_ArbitraryFloatATanINTEL{{.*}}(i44 {{[%a-z0-9.]+}}, i32 31, i32 31, i32 0, i32 2, i32 1) @@ -344,8 +345,8 @@ void ap_float_atan() { template void ap_float_atapin() { - ap_int<1 + EA + MA> A; - ap_int<1 + Eout + Mout> atanpi_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + Eout + Mout> atanpi_res = __spirv_ArbitraryFloatATanPiINTEL<1 + EA + MA, 1 + Eout + Mout>( A, MA, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i34 @_Z{{[0-9]+}}__spirv_ArbitraryFloatATanPiINTEL{{.*}}(i40 {{[%a-z0-9.]+}}, i32 38, i32 32, i32 0, i32 2, i32 1) @@ -353,19 +354,20 @@ void ap_float_atapin() { template void ap_float_atan2() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; - ap_int<1 + Eout + Mout> atan2_res = - __spirv_ArbitraryFloatATan2INTEL<1 + EA + MA, 1 + EB + MB, 1 + Eout + Mout>( + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + Eout + Mout> atan2_res = + __spirv_ArbitraryFloatATan2INTEL<1 + EA + MA, 1 + EB + MB, + 1 + Eout + Mout>( A, MA, B, MB, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i27 @_Z{{[0-9]+}}__spirv_ArbitraryFloatATan2INTEL{{.*}}(i24 signext {{[%a-z0-9.]+}}, i32 16, i25 signext {{[%a-z0-9.]+}}, i32 17, i32 18, i32 0, i32 2, i32 1) } template void ap_float_pow() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; - ap_int<1 + Eout + Mout> pow_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + Eout + Mout> pow_res = __spirv_ArbitraryFloatPowINTEL<1 + EA + MA, 1 + EB + MB, 1 + Eout + Mout>( A, MA, B, MB, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func signext i21 @_Z{{[0-9]+}}__spirv_ArbitraryFloatPowINTEL{{.*}}(i17 signext {{[%a-z0-9.]+}}, i32 8, i19 signext {{[%a-z0-9.]+}}, i32 9, i32 10, i32 0, i32 2, i32 1) @@ -373,19 +375,20 @@ void ap_float_pow() { template void ap_float_powr() { - ap_int<1 + EA + MA> A; - ap_int<1 + EB + MB> B; - ap_int<1 + Eout + Mout> powr_res = - __spirv_ArbitraryFloatPowRINTEL<1 + EA + MA, 1 + EB + MB, 1 + Eout + Mout>( + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int<1 + EB + MB> B; + sycl::detail::ap_int<1 + Eout + Mout> powr_res = + __spirv_ArbitraryFloatPowRINTEL<1 + EA + MA, 1 + EB + MB, + 1 + Eout + Mout>( A, MA, B, MB, Mout, Subnorm, RndMode, RndAcc); // CHECK: call spir_func i56 @_Z{{[0-9]+}}__spirv_ArbitraryFloatPowRINTEL{{.*}}(i54 {{[%a-z0-9.]+}}, i32 35, i55 {{[%a-z0-9.]+}}, i32 35, i32 35, i32 0, i32 2, i32 1) } template void ap_float_pown() { - ap_int<1 + EA + MA> A; - ap_int B; - ap_int<1 + Eout + Mout> pown_res = + sycl::detail::ap_int<1 + EA + MA> A; + sycl::detail::ap_int 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)