[SYCL] Adding support for missing math ops (intel#14132)

[SYCL] Adding support for missing math ops: - truncf - sinpif - rsqrtf - exp10f
ianayl · Jun 13, 2024 · fd07268 · fd07268
1 parent a48a126
commit fd07268
Show file tree

Hide file tree

Showing 7 changed files with 92 additions and 12 deletions.
diff --git a/libdevice/cmath_wrapper.cpp b/libdevice/cmath_wrapper.cpp
@@ -39,6 +39,18 @@ extern "C" SYCL_EXTERNAL float __devicelib_fminf(float, float);
 DEVICE_EXTERN_C_INLINE
 float fminf(float x, float y) { return __devicelib_fminf(x, y); }
 
+DEVICE_EXTERN_C_INLINE
+float truncf(float x) { return __devicelib_truncf(x); }
+
+DEVICE_EXTERN_C_INLINE
+float sinpif(float x) { return __devicelib_sinpif(x); }
+
+DEVICE_EXTERN_C_INLINE
+float rsqrtf(float x) { return __devicelib_rsqrtf(x); }
+
+DEVICE_EXTERN_C_INLINE
+float exp10f(float x) { return __devicelib_exp10f(x); }
+
 DEVICE_EXTERN_C_INLINE
 div_t div(int x, int y) { return __devicelib_div(x, y); }
 

diff --git a/libdevice/cmath_wrapper_fp64.cpp b/libdevice/cmath_wrapper_fp64.cpp
@@ -36,6 +36,18 @@ extern "C" SYCL_EXTERNAL double __devicelib_fmin(double, double);
 DEVICE_EXTERN_C_INLINE
 double fmin(double x, double y) { return __devicelib_fmin(x, y); }
 
+DEVICE_EXTERN_C_INLINE
+double trunc(double x) { return __devicelib_trunc(x); }
+
+DEVICE_EXTERN_C_INLINE
+double sinpi(double x) { return __devicelib_sinpi(x); }
+
+DEVICE_EXTERN_C_INLINE
+double rsqrt(double x) { return __devicelib_rsqrt(x); }
+
+DEVICE_EXTERN_C_INLINE
+double exp10(double x) { return __devicelib_exp10(x); }
+
 DEVICE_EXTERN_C_INLINE
 double log(double x) { return __devicelib_log(x); }
 

diff --git a/libdevice/device_math.h b/libdevice/device_math.h
@@ -76,6 +76,30 @@ float __devicelib_fminf(float x, float y);
 DEVICE_EXTERN_C
 double __devicelib_fmin(double x, double y);
 
+DEVICE_EXTERN_C
+float __devicelib_truncf(float x);
+
+DEVICE_EXTERN_C
+double __devicelib_trunc(double x);
+
+DEVICE_EXTERN_C
+double __devicelib_sinpi(double x);
+
+DEVICE_EXTERN_C
+float __devicelib_sinpif(float x);
+
+DEVICE_EXTERN_C
+double __devicelib_rsqrt(double x);
+
+DEVICE_EXTERN_C
+float __devicelib_rsqrtf(float x);
+
+DEVICE_EXTERN_C
+double __devicelib_exp10(double x);
+
+DEVICE_EXTERN_C
+float __devicelib_exp10f(float x);
+
 DEVICE_EXTERN_C
 div_t __devicelib_div(int x, int y);
 

diff --git a/libdevice/fallback-cmath-fp64.cpp b/libdevice/fallback-cmath-fp64.cpp
@@ -35,6 +35,18 @@ double __devicelib_fmax(double x, double y) { return __spirv_ocl_fmax(x, y); }
 DEVICE_EXTERN_C_INLINE
 double __devicelib_fmin(double x, double y) { return __spirv_ocl_fmin(x, y); }
 
+DEVICE_EXTERN_C_INLINE
+double __devicelib_trunc(double x) { return __spirv_ocl_trunc(x); }
+
+DEVICE_EXTERN_C_INLINE
+double __devicelib_sinpi(double x) { return __spirv_ocl_sinpi(x); }
+
+DEVICE_EXTERN_C_INLINE
+double __devicelib_rsqrt(double x) { return __spirv_ocl_rsqrt(x); }
+
+DEVICE_EXTERN_C_INLINE
+double __devicelib_exp10(double x) { return __spirv_ocl_exp10(x); }
+
 DEVICE_EXTERN_C_INLINE
 double __devicelib_log(double x) { return __spirv_ocl_log(x); }
 

diff --git a/libdevice/fallback-cmath.cpp b/libdevice/fallback-cmath.cpp
@@ -45,6 +45,18 @@ float __devicelib_fmaxf(float x, float y) { return __spirv_ocl_fmax(x, y); }
 DEVICE_EXTERN_C_INLINE
 float __devicelib_fminf(float x, float y) { return __spirv_ocl_fmin(x, y); }
 
+DEVICE_EXTERN_C_INLINE
+float __devicelib_truncf(float x) { return __spirv_ocl_trunc(x); }
+
+DEVICE_EXTERN_C_INLINE
+float __devicelib_sinpif(float x) { return __spirv_ocl_sinpi(x); }
+
+DEVICE_EXTERN_C_INLINE
+float __devicelib_rsqrtf(float x) { return __spirv_ocl_rsqrt(x); }
+
+DEVICE_EXTERN_C_INLINE
+float __devicelib_exp10f(float x) { return __spirv_ocl_exp10(x); }
+
 DEVICE_EXTERN_C_INLINE
 div_t __devicelib_div(int x, int y) { return {x / y, x % y}; }
 

diff --git a/sycl/test-e2e/DeviceLib/cmath_fp64_test.cpp b/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 69
+#define TEST_NUM 73
 
-double ref[TEST_NUM] = {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 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};
 
 double refIptr = 1;
 
@@ -61,6 +61,10 @@ template <class T> void device_cmath_test(s::queue &deviceQueue) {
         T minus_infinity = -INFINITY;
         double subnormal;
         *((uint64_t *)&subnormal) = 0xFFFFFFFFFFFFFULL;
+        res_access[i++] = sycl::exp10(2.0);
+        res_access[i++] = sycl::rsqrt(4.0);
+        res_access[i++] = std::trunc(1.3);
+        res_access[i++] = sycl::sinpi(0.0);
         res_access[i++] = sycl::cospi(0.5);
         res_access[i++] = std::copysign(2, -1);
         res_access[i++] = std::fmin(2, 1);

diff --git a/sycl/test-e2e/DeviceLib/cmath_test.cpp b/sycl/test-e2e/DeviceLib/cmath_test.cpp
@@ -22,13 +22,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 66
+#define TEST_NUM 70
 
-float ref[TEST_NUM] = {0,   -2, 1,   2, 1, 1, 0, 1, 1, 0, 0, 0,   0,   0,
-                       1,   1,  0.5, 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};
+float ref[TEST_NUM] = {100, 0.5, 1.0, 0,   0,   -2, 1,   2, 1, 1, 0, 1, 1, 0,
+                       0,   0,   0,   0,   1,   1,  0.5, 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};
 
 float refIptr = 1;
 
@@ -58,6 +58,10 @@ template <class T> void device_cmath_test_1(s::queue &deviceQueue) {
         float subnormal;
         *((uint32_t *)&subnormal) = 0x7FFFFF;
 
+        res_access[i++] = sycl::exp10(2.0f);
+        res_access[i++] = sycl::rsqrt(4.0f);
+        res_access[i++] = std::trunc(1.2f);
+        res_access[i++] = sycl::sinpi(0.0f);
         res_access[i++] = sycl::cospi(0.5f);
         res_access[i++] = std::copysign(2.0f, -10.0f);
         res_access[i++] = sycl::min(2.0f, 1.0f);