[SYCL] Add support for scalbln (#14401)

Add support for missing scalbln function

Author: MaryaSharf

libdevice/cmath_wrapper.cpp

@@ -69,6 +69,9 @@ float floorf(float x) { return __devicelib_floorf(x); }
 DEVICE_EXTERN_C_INLINE
 float scalbnf(float x, int n) { return __devicelib_scalbnf(x, n); }
 
+DEVICE_EXTERN_C_INLINE
+float scalblnf(float x, long int n) { return __devicelib_scalblnf(x, n); }
+
 DEVICE_EXTERN_C_INLINE
 float logf(float x) { return __devicelib_logf(x); }
 

libdevice/cmath_wrapper_fp64.cpp

@@ -25,6 +25,9 @@ double ceil(double x) { return __devicelib_ceil(x); }
 DEVICE_EXTERN_C_INLINE
 double copysign(double x, double y) { return __devicelib_copysign(x, y); }
 
+DEVICE_EXTERN_C_INLINE
+double scalbln(double x, long y) { return __devicelib_scalbln(x, y); }
+
 DEVICE_EXTERN_C_INLINE
 double cospi(double x) { return __devicelib_cospi(x); }
 

libdevice/device_math.h

@@ -64,6 +64,12 @@ float __devicelib_cospif(float x);
 DEVICE_EXTERN_C
 double __devicelib_cospi(double x);
 
+DEVICE_EXTERN_C
+float __devicelib_scalblnf(float x, long int y);
+
+DEVICE_EXTERN_C
+double __devicelib_scalbln(double x, long int y);
+
 DEVICE_EXTERN_C
 float __devicelib_fmaxf(float x, float y);
 

libdevice/fallback-cmath-fp64.cpp

@@ -29,6 +29,11 @@ double __devicelib_copysign(double x, double y) {
 DEVICE_EXTERN_C_INLINE
 double __devicelib_cospi(double x) { return __spirv_ocl_cospi(x); }
 
+DEVICE_EXTERN_C_INLINE
+double __devicelib_scalbln(double x, long int y) {
+  return __spirv_ocl_ldexp(x, (int)y);
+}
+
 DEVICE_EXTERN_C_INLINE
 double __devicelib_fmax(double x, double y) { return __spirv_ocl_fmax(x, y); }
 

libdevice/fallback-cmath.cpp

@@ -39,6 +39,11 @@ float __devicelib_copysignf(float x, float y) {
 DEVICE_EXTERN_C_INLINE
 float __devicelib_cospif(float x) { return __spirv_ocl_cospi(x); }
 
+DEVICE_EXTERN_C_INLINE
+float __devicelib_scalblnf(float x, long int y) {
+  return __spirv_ocl_ldexp(x, (int)y);
+}
+
 DEVICE_EXTERN_C_INLINE
 float __devicelib_fmaxf(float x, float y) { return __spirv_ocl_fmax(x, y); }
 

sycl/test-e2e/DeviceLib/cmath_fp64_test.cpp

@@ -20,13 +20,13 @@ namespace s = sycl;
 constexpr s::access::mode sycl_read = s::access::mode::read;
 constexpr s::access::mode sycl_write = s::access::mode::write;
 
-#define TEST_NUM 73
+#define TEST_NUM 74
 
 double ref[TEST_NUM] = {
-    100, 0.5, 1.0, 0, 0, -2, 1, 2, 1,   1,   1,   0,   1, 1, 0, 0, 0, 0, 0,
-    1,   1,   0.5, 0, 2, 0,  0, 1, 0,   2,   0,   0,   0, 0, 0, 1, 0, 1, 2,
-    0,   1,   2,   5, 0, 0,  0, 0, 0.5, 0.5, NAN, NAN, 2, 0, 0, 0, 0, 0, 0,
-    0,   0,   0,   0, 0, 0,  0, 0, 0,   0,   0,   0,   0, 0, 0, 0};
+    6, 100, 0.5, 1.0, 0, 0, -2, 1, 2, 1,   1,   1,   0,   1, 1, 0, 0, 0, 0,
+    0, 1,   1,   0.5, 0, 2, 0,  0, 1, 0,   2,   0,   0,   0, 0, 0, 1, 0, 1,
+    2, 0,   1,   2,   5, 0, 0,  0, 0, 0.5, 0.5, NAN, NAN, 2, 0, 0, 0, 0, 0,
+    0, 0,   0,   0,   0, 0, 0,  0, 0, 0,   0,   0,   0,   0, 0, 0, 0};
 
 double refIptr = 1;
 
@@ -61,6 +61,7 @@ template <class T> void device_cmath_test(s::queue &deviceQueue) {
         T minus_infinity = -INFINITY;
         double subnormal;
         *((uint64_t *)&subnormal) = 0xFFFFFFFFFFFFFULL;
+        res_access[i++] = std::scalbln(1.5, 2);
         res_access[i++] = sycl::exp10(2.0);
         res_access[i++] = sycl::rsqrt(4.0);
         res_access[i++] = std::trunc(1.3);

sycl/test-e2e/DeviceLib/cmath_test.cpp

@@ -15,6 +15,7 @@
 #include <cmath>
 #include <cstdint>
 #include <iostream>
+#include <std/experimental/simd.hpp>
 #include <sycl/detail/core.hpp>
 #include <sycl/usm.hpp>
 
@@ -154,15 +155,15 @@ template <class T> void device_cmath_test_1(s::queue &deviceQueue) {
   assert(quo == 0);
 }
 
-// MSVC implements std::ldexp<float>, std::fabs<float> and std::frexp<float> by
-// invoking the 'double' version of corresponding C math functions(ldexp, fabs
-// and frexp). Those functions can only work on Windows with fp64 extension
-// support from underlying device.
+// MSVC implements std::scalbln<float>, std::ldexp<float>, std::fabs<float> and
+// std::frexp<float> by invoking the 'double' version of corresponding C math
+// functions(scalbln, ldexp, fabs and frexp). Those functions can only work on
+// Windows with fp64 extension support from underlying device.
 #ifndef _WIN32
 template <class T> void device_cmath_test_2(s::queue &deviceQueue) {
-  constexpr int NumOfTestItems = 3;
+  constexpr int NumOfTestItems = 4;
   T result[NumOfTestItems] = {-1};
-  T ref[NumOfTestItems] = {0, 2, 1};
+  T ref[NumOfTestItems] = {6, 0, 2, 1};
   // Variable exponent is an integer value to store the exponent in frexp
   // function
   int exponent = -1;
@@ -175,6 +176,7 @@ template <class T> void device_cmath_test_2(s::queue &deviceQueue) {
       auto exp_access = buffer2.template get_access<sycl_write>(cgh);
       cgh.single_task<class DeviceMathTest2>([=]() {
         int i = 0;
+        res_access[i++] = std::scalbln(1.5f, 2);
         res_access[i++] = std::frexp(0.0f, &exp_access[0]);
         res_access[i++] = std::ldexp(1.0f, 1);
         res_access[i++] = std::fabs(-1.0f);