Skip to content

Commit 51b6d1c

Browse files
MaryaSharfMaryaSharf
committed
Adding support for missing math ops
- truncf - sinpif - rsqrtf - exp10f
1 parent 6ecce4f commit 51b6d1c

File tree

7 files changed

+92
-12
lines changed

7 files changed

+92
-12
lines changed

‎libdevice/cmath_wrapper.cpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -39,6 +39,18 @@ extern "C" SYCL_EXTERNAL float __devicelib_fminf(float, float);
3939
DEVICE_EXTERN_C_INLINE
4040
float fminf(float x, float y) { return __devicelib_fminf(x, y); }
4141

42+
DEVICE_EXTERN_C_INLINE
43+
float truncf(float x) { return __devicelib_truncf(x); }
44+
45+
DEVICE_EXTERN_C_INLINE
46+
float sinpif(float x) { return __devicelib_sinpif(x); }
47+
48+
DEVICE_EXTERN_C_INLINE
49+
float rsqrtf(float x) { return __devicelib_rsqrtf(x); }
50+
51+
DEVICE_EXTERN_C_INLINE
52+
float exp10f(float x) { return __devicelib_exp10f(x); }
53+
4254
DEVICE_EXTERN_C_INLINE
4355
div_t div(int x, int y) { return __devicelib_div(x, y); }
4456

‎libdevice/cmath_wrapper_fp64.cpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -36,6 +36,18 @@ extern "C" SYCL_EXTERNAL double __devicelib_fmin(double, double);
3636
DEVICE_EXTERN_C_INLINE
3737
double fmin(double x, double y) { return __devicelib_fmin(x, y); }
3838

39+
DEVICE_EXTERN_C_INLINE
40+
double trunc(double x) { return __devicelib_trunc(x); }
41+
42+
DEVICE_EXTERN_C_INLINE
43+
double sinpi(double x) { return __devicelib_sinpi(x); }
44+
45+
DEVICE_EXTERN_C_INLINE
46+
double rsqrt(double x) { return __devicelib_rsqrt(x); }
47+
48+
DEVICE_EXTERN_C_INLINE
49+
double exp10(double x) { return __devicelib_exp10(x); }
50+
3951
DEVICE_EXTERN_C_INLINE
4052
double log(double x) { return __devicelib_log(x); }
4153

‎libdevice/device_math.h

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -76,6 +76,30 @@ float __devicelib_fminf(float x, float y);
7676
DEVICE_EXTERN_C
7777
double __devicelib_fmin(double x, double y);
7878

79+
DEVICE_EXTERN_C
80+
float __devicelib_truncf(float x);
81+
82+
DEVICE_EXTERN_C
83+
double __devicelib_trunc(double x);
84+
85+
DEVICE_EXTERN_C
86+
double __devicelib_sinpi(double x);
87+
88+
DEVICE_EXTERN_C
89+
float __devicelib_sinpif(float x);
90+
91+
DEVICE_EXTERN_C
92+
double __devicelib_rsqrt(double x);
93+
94+
DEVICE_EXTERN_C
95+
float __devicelib_rsqrtf(float x);
96+
97+
DEVICE_EXTERN_C
98+
double __devicelib_exp10(double x);
99+
100+
DEVICE_EXTERN_C
101+
float __devicelib_exp10f(float x);
102+
79103
DEVICE_EXTERN_C
80104
div_t __devicelib_div(int x, int y);
81105

‎libdevice/fallback-cmath-fp64.cpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -35,6 +35,18 @@ double __devicelib_fmax(double x, double y) { return __spirv_ocl_fmax(x, y); }
3535
DEVICE_EXTERN_C_INLINE
3636
double __devicelib_fmin(double x, double y) { return __spirv_ocl_fmin(x, y); }
3737

38+
DEVICE_EXTERN_C_INLINE
39+
double __devicelib_trunc(double x) { return __spirv_ocl_trunc(x); }
40+
41+
DEVICE_EXTERN_C_INLINE
42+
double __devicelib_sinpi(double x) { return __spirv_ocl_sinpi(x); }
43+
44+
DEVICE_EXTERN_C_INLINE
45+
double __devicelib_rsqrt(double x) { return __spirv_ocl_rsqrt(x); }
46+
47+
DEVICE_EXTERN_C_INLINE
48+
double __devicelib_exp10(double x) { return __spirv_ocl_exp10(x); }
49+
3850
DEVICE_EXTERN_C_INLINE
3951
double __devicelib_log(double x) { return __spirv_ocl_log(x); }
4052

‎libdevice/fallback-cmath.cpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -45,6 +45,18 @@ float __devicelib_fmaxf(float x, float y) { return __spirv_ocl_fmax(x, y); }
4545
DEVICE_EXTERN_C_INLINE
4646
float __devicelib_fminf(float x, float y) { return __spirv_ocl_fmin(x, y); }
4747

48+
DEVICE_EXTERN_C_INLINE
49+
float __devicelib_truncf(float x) { return __spirv_ocl_trunc(x); }
50+
51+
DEVICE_EXTERN_C_INLINE
52+
float __devicelib_sinpif(float x) { return __spirv_ocl_sinpi(x); }
53+
54+
DEVICE_EXTERN_C_INLINE
55+
float __devicelib_rsqrtf(float x) { return __spirv_ocl_rsqrt(x); }
56+
57+
DEVICE_EXTERN_C_INLINE
58+
float __devicelib_exp10f(float x) { return __spirv_ocl_exp10(x); }
59+
4860
DEVICE_EXTERN_C_INLINE
4961
div_t __devicelib_div(int x, int y) { return {x / y, x % y}; }
5062

‎sycl/test-e2e/DeviceLib/cmath_fp64_test.cpp

Lines changed: 10 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -20,13 +20,13 @@ namespace s = sycl;
2020
constexpr s::access::mode sycl_read = s::access::mode::read;
2121
constexpr s::access::mode sycl_write = s::access::mode::write;
2222

23-
#define TEST_NUM 69
23+
#define TEST_NUM 73
2424

25-
double ref[TEST_NUM] = {0, -2, 1, 2, 1, 1, 1, 0, 1, 1, 0, 0, 0, 0,
26-
0, 1, 1, 0.5, 0, 2, 0, 0, 1, 0, 2, 0, 0, 0,
27-
0, 0, 1, 0, 1, 2, 0, 1, 2, 5, 0, 0, 0, 0,
28-
0.5, 0.5, NAN, NAN, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0,
29-
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
25+
double ref[TEST_NUM] = {
26+
100, 0.5, 1.0, 0, 0, -2, 1, 2, 1, 1, 1, 0, 1, 1, 0, 0, 0, 0, 0,
27+
1, 1, 0.5, 0, 2, 0, 0, 1, 0, 2, 0, 0, 0, 0, 0, 1, 0, 1, 2,
28+
0, 1, 2, 5, 0, 0, 0, 0, 0.5, 0.5, NAN, NAN, 2, 0, 0, 0, 0, 0, 0,
29+
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
3030

3131
double refIptr = 1;
3232

@@ -61,6 +61,10 @@ template <class T> void device_cmath_test(s::queue &deviceQueue) {
6161
T minus_infinity = -INFINITY;
6262
double subnormal;
6363
*((uint64_t *)&subnormal) = 0xFFFFFFFFFFFFFULL;
64+
res_access[i++] = sycl::exp10(2.0);
65+
res_access[i++] = sycl::rsqrt(4.0);
66+
res_access[i++] = std::trunc(1.3);
67+
res_access[i++] = sycl::sinpi(0.0);
6468
res_access[i++] = sycl::cospi(0.5);
6569
res_access[i++] = std::copysign(2, -1);
6670
res_access[i++] = std::fmin(2, 1);

‎sycl/test-e2e/DeviceLib/cmath_test.cpp

Lines changed: 10 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -22,13 +22,13 @@ namespace s = sycl;
2222
constexpr s::access::mode sycl_read = s::access::mode::read;
2323
constexpr s::access::mode sycl_write = s::access::mode::write;
2424

25-
#define TEST_NUM 66
25+
#define TEST_NUM 70
2626

27-
float ref[TEST_NUM] = {0, -2, 1, 2, 1, 1, 0, 1, 1, 0, 0, 0, 0, 0,
28-
1, 1, 0.5, 0, 0, 1, 0, 2, 0, 0, 0, 0, 0, 1,
29-
0, 1, 2, 0, 1, 2, 5, 0, 0, 0, 0, 0.5, 0.5, NAN,
30-
NAN, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
31-
0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
27+
float ref[TEST_NUM] = {100, 0.5, 1.0, 0, 0, -2, 1, 2, 1, 1, 0, 1, 1, 0,
28+
0, 0, 0, 0, 1, 1, 0.5, 0, 0, 1, 0, 2, 0, 0,
29+
0, 0, 0, 1, 0, 1, 2, 0, 1, 2, 5, 0, 0, 0,
30+
0, 0.5, 0.5, NAN, NAN, 2, 0, 0, 0, 0, 0, 0, 0, 0,
31+
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
3232

3333
float refIptr = 1;
3434

@@ -58,6 +58,10 @@ template <class T> void device_cmath_test_1(s::queue &deviceQueue) {
5858
float subnormal;
5959
*((uint32_t *)&subnormal) = 0x7FFFFF;
6060

61+
res_access[i++] = sycl::exp10(2.0f);
62+
res_access[i++] = sycl::rsqrt(4.0f);
63+
res_access[i++] = std::trunc(1.2f);
64+
res_access[i++] = sycl::sinpi(0.0f);
6165
res_access[i++] = sycl::cospi(0.5f);
6266
res_access[i++] = std::copysign(2.0f, -10.0f);
6367
res_access[i++] = sycl::min(2.0f, 1.0f);

0 commit comments

Comments
 (0)