Skip to content

[SYCL] Adding support for missing math ops #14132

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Jun 13, 2024
Merged

[SYCL] Adding support for missing math ops #14132

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 12 additions & 0 deletions libdevice/cmath_wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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); }

Expand Down
12 changes: 12 additions & 0 deletions libdevice/cmath_wrapper_fp64.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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); }

Expand Down
24 changes: 24 additions & 0 deletions libdevice/device_math.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand Down
12 changes: 12 additions & 0 deletions libdevice/fallback-cmath-fp64.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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); }

Expand Down
12 changes: 12 additions & 0 deletions libdevice/fallback-cmath.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}; }

Expand Down
16 changes: 10 additions & 6 deletions sycl/test-e2e/DeviceLib/cmath_fp64_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;

Expand Down Expand Up @@ -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);
Expand Down
16 changes: 10 additions & 6 deletions sycl/test-e2e/DeviceLib/cmath_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;

Expand Down Expand Up @@ -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);
Expand Down
Loading