Skip to content

Commit da379ec

Browse files
MaryaSharfMaryaSharf
authored
[SYCL] Add support for multiple missing math ops (intel#13714)
Add support for multiple missing math ops; - ceilf - copysignf - cospif - fmaxf - fminf
1 parent 84bae21 commit da379ec

File tree

7 files changed

+122
-10
lines changed

7 files changed

+122
-10
lines changed

‎libdevice/cmath_wrapper.cpp‎

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,23 @@ long long int llabs(long long int x) { return __devicelib_llabs(x); }
2222
DEVICE_EXTERN_C_INLINE
2323
float fabsf(float x) { return __devicelib_fabsf(x); }
2424

25+
DEVICE_EXTERN_C_INLINE
26+
float ceilf(float x) { return __devicelib_ceilf(x); }
27+
28+
DEVICE_EXTERN_C_INLINE
29+
float copysignf(float x, float y) { return __devicelib_copysignf(x, y); }
30+
31+
DEVICE_EXTERN_C_INLINE
32+
float cospif(float x) { return __devicelib_cospif(x); }
33+
34+
extern "C" SYCL_EXTERNAL float __devicelib_fmaxf(float, float);
35+
DEVICE_EXTERN_C_INLINE
36+
float fmaxf(float x, float y) { return __devicelib_fmaxf(x, y); }
37+
38+
extern "C" SYCL_EXTERNAL float __devicelib_fminf(float, float);
39+
DEVICE_EXTERN_C_INLINE
40+
float fminf(float x, float y) { return __devicelib_fminf(x, y); }
41+
2542
DEVICE_EXTERN_C_INLINE
2643
div_t div(int x, int y) { return __devicelib_div(x, y); }
2744

‎libdevice/cmath_wrapper_fp64.cpp‎

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,23 @@
1919
DEVICE_EXTERN_C_INLINE
2020
double fabs(double x) { return __devicelib_fabs(x); }
2121

22+
DEVICE_EXTERN_C_INLINE
23+
double ceil(double x) { return __devicelib_ceil(x); }
24+
25+
DEVICE_EXTERN_C_INLINE
26+
double copysign(double x, double y) { return __devicelib_copysign(x, y); }
27+
28+
DEVICE_EXTERN_C_INLINE
29+
double cospi(double x) { return __devicelib_cospi(x); }
30+
31+
extern "C" SYCL_EXTERNAL double __devicelib_fmax(double, double);
32+
DEVICE_EXTERN_C_INLINE
33+
double fmax(double x, double y) { return __devicelib_fmax(x, y); }
34+
35+
extern "C" SYCL_EXTERNAL double __devicelib_fmin(double, double);
36+
DEVICE_EXTERN_C_INLINE
37+
double fmin(double x, double y) { return __devicelib_fmin(x, y); }
38+
2239
DEVICE_EXTERN_C_INLINE
2340
double log(double x) { return __devicelib_log(x); }
2441

‎libdevice/device_math.h‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -46,6 +46,36 @@ float __devicelib_fabsf(float x);
4646
DEVICE_EXTERN_C
4747
double __devicelib_fabs(double x);
4848

49+
DEVICE_EXTERN_C
50+
float __devicelib_ceilf(float x);
51+
52+
DEVICE_EXTERN_C
53+
double __devicelib_ceil(double x);
54+
55+
DEVICE_EXTERN_C
56+
float __devicelib_copysignf(float x, float y);
57+
58+
DEVICE_EXTERN_C
59+
double __devicelib_copysign(double x, double y);
60+
61+
DEVICE_EXTERN_C
62+
float __devicelib_cospif(float x);
63+
64+
DEVICE_EXTERN_C
65+
double __devicelib_cospi(double x);
66+
67+
DEVICE_EXTERN_C
68+
float __devicelib_fmaxf(float x, float y);
69+
70+
DEVICE_EXTERN_C
71+
double __devicelib_fmax(double x, double y);
72+
73+
DEVICE_EXTERN_C
74+
float __devicelib_fminf(float x, float y);
75+
76+
DEVICE_EXTERN_C
77+
double __devicelib_fmin(double x, double y);
78+
4979
DEVICE_EXTERN_C
5080
div_t __devicelib_div(int x, int y);
5181

‎libdevice/fallback-cmath-fp64.cpp‎

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,23 @@
1818
DEVICE_EXTERN_C_INLINE
1919
double __devicelib_fabs(double x) { return x < 0 ? -x : x; }
2020

21+
DEVICE_EXTERN_C_INLINE
22+
double __devicelib_ceil(double x) { return __spirv_ocl_ceil(x); }
23+
24+
DEVICE_EXTERN_C_INLINE
25+
double __devicelib_copysign(double x, double y) {
26+
return __spirv_ocl_copysign(x, y);
27+
}
28+
29+
DEVICE_EXTERN_C_INLINE
30+
double __devicelib_cospi(double x) { return __spirv_ocl_cospi(x); }
31+
32+
DEVICE_EXTERN_C_INLINE
33+
double __devicelib_fmax(double x, double y) { return __spirv_ocl_fmax(x, y); }
34+
35+
DEVICE_EXTERN_C_INLINE
36+
double __devicelib_fmin(double x, double y) { return __spirv_ocl_fmin(x, y); }
37+
2138
DEVICE_EXTERN_C_INLINE
2239
double __devicelib_log(double x) { return __spirv_ocl_log(x); }
2340

‎libdevice/fallback-cmath.cpp‎

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -28,6 +28,23 @@ long long int __devicelib_llabs(long long int x) { return x < 0 ? -x : x; }
2828
DEVICE_EXTERN_C_INLINE
2929
float __devicelib_fabsf(float x) { return x < 0 ? -x : x; }
3030

31+
DEVICE_EXTERN_C_INLINE
32+
float __devicelib_ceilf(float x) { return __spirv_ocl_ceil(x); }
33+
34+
DEVICE_EXTERN_C_INLINE
35+
float __devicelib_copysignf(float x, float y) {
36+
return __spirv_ocl_copysign(x, y);
37+
}
38+
39+
DEVICE_EXTERN_C_INLINE
40+
float __devicelib_cospif(float x) { return __spirv_ocl_cospi(x); }
41+
42+
DEVICE_EXTERN_C_INLINE
43+
float __devicelib_fmaxf(float x, float y) { return __spirv_ocl_fmax(x, y); }
44+
45+
DEVICE_EXTERN_C_INLINE
46+
float __devicelib_fminf(float x, float y) { return __spirv_ocl_fmin(x, y); }
47+
3148
DEVICE_EXTERN_C_INLINE
3249
div_t __devicelib_div(int x, int y) { return {x / y, x % y}; }
3350

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

Lines changed: 12 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -14,17 +14,19 @@
1414
#include <cstdint>
1515
#include <iostream>
1616
#include <sycl/detail/core.hpp>
17+
#include <sycl/usm.hpp>
1718

1819
namespace s = sycl;
1920
constexpr s::access::mode sycl_read = s::access::mode::read;
2021
constexpr s::access::mode sycl_write = s::access::mode::write;
2122

22-
#define TEST_NUM 64
23+
#define TEST_NUM 69
2324

24-
double ref[TEST_NUM] = {
25-
1, 1, 0, 1, 1, 0, 0, 0, 0, 0, 1, 1, 0.5, 0, 2, 0, 0, 1, 0, 2, 0, 0,
26-
0, 0, 0, 1, 0, 1, 2, 0, 1, 2, 5, 0, 0, 0, 0, 0.5, 0.5, NAN, NAN, 2, 0, 0,
27-
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
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};
2830

2931
double refIptr = 1;
3032

@@ -59,7 +61,12 @@ template <class T> void device_cmath_test(s::queue &deviceQueue) {
5961
T minus_infinity = -INFINITY;
6062
double subnormal;
6163
*((uint64_t *)&subnormal) = 0xFFFFFFFFFFFFFULL;
64+
res_access[i++] = sycl::cospi(0.5);
65+
res_access[i++] = std::copysign(2, -1);
66+
res_access[i++] = std::fmin(2, 1);
67+
res_access[i++] = std::fmax(2, 1);
6268
res_access[i++] = std::fabs(-1.0);
69+
res_access[i++] = std::ceil(0.1);
6370
res_access[i++] = std::cos(0.0);
6471
res_access[i++] = std::sin(0.0);
6572
res_access[i++] = std::round(1.0);

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

Lines changed: 12 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -16,17 +16,19 @@
1616
#include <cstdint>
1717
#include <iostream>
1818
#include <sycl/detail/core.hpp>
19+
#include <sycl/usm.hpp>
1920

2021
namespace s = sycl;
2122
constexpr s::access::mode sycl_read = s::access::mode::read;
2223
constexpr s::access::mode sycl_write = s::access::mode::write;
2324

24-
#define TEST_NUM 61
25+
#define TEST_NUM 66
2526

26-
float ref[TEST_NUM] = {1, 0, 1, 1, 0, 0, 0, 0, 0, 1, 1, 0.5, 0, 0, 1, 0,
27-
2, 0, 0, 0, 0, 0, 1, 0, 1, 2, 0, 1, 2, 5, 0, 0,
28-
0, 0, 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};
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};
3032

3133
float refIptr = 1;
3234

@@ -56,6 +58,11 @@ template <class T> void device_cmath_test_1(s::queue &deviceQueue) {
5658
float subnormal;
5759
*((uint32_t *)&subnormal) = 0x7FFFFF;
5860

61+
res_access[i++] = sycl::cospi(0.5f);
62+
res_access[i++] = std::copysign(2.0f, -10.0f);
63+
res_access[i++] = sycl::min(2.0f, 1.0f);
64+
res_access[i++] = sycl::max(2.0f, 1.0f);
65+
res_access[i++] = sycl::ceil(0.1f);
5966
res_access[i++] = std::cos(0.0f);
6067
res_access[i++] = std::sin(0.0f);
6168
res_access[i++] = std::round(1.0f);

0 commit comments

Comments
 (0)