Skip to content

Commit 7abd9d5

Browse files
jinge90jinge90
authored
[SYCL] Add device libraries for C math and complex (#975)
Add device libraries for C math and complex functions based on SYCL compiler's device library infrastructure. Fallback implementation for math and complex functions are also provided. The fallback implementation for C math functions are based on `__spirv_ocl*` math built-ins and the fallback implementation for C99 complex functions are based on algorithm used in libc++ std::complex. All functions supported are listed in sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst. Signed-off-by: gejin <[email protected]>
1 parent ffcad03 commit 7abd9d5

27 files changed

+4637
-35
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 26 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -138,48 +138,52 @@ static bool IsSyclMathFunc(unsigned BuiltinID) {
138138
case Builtin::BI__builtin_truncl:
139139
case Builtin::BIlroundl:
140140
case Builtin::BI__builtin_lroundl:
141+
case Builtin::BIceil:
142+
case Builtin::BI__builtin_ceil:
143+
case Builtin::BIcopysign:
144+
case Builtin::BI__builtin_copysign:
145+
case Builtin::BIfabs:
146+
case Builtin::BI__builtin_fabs:
147+
case Builtin::BIfloor:
148+
case Builtin::BI__builtin_floor:
149+
case Builtin::BIfmax:
150+
case Builtin::BI__builtin_fmax:
151+
case Builtin::BIfmin:
152+
case Builtin::BI__builtin_fmin:
153+
case Builtin::BInearbyint:
154+
case Builtin::BI__builtin_nearbyint:
155+
case Builtin::BIrint:
156+
case Builtin::BI__builtin_rint:
157+
case Builtin::BIround:
158+
case Builtin::BI__builtin_round:
159+
case Builtin::BItrunc:
160+
case Builtin::BI__builtin_trunc:
141161
case Builtin::BIceilf:
142162
case Builtin::BI__builtin_ceilf:
143163
case Builtin::BIcopysignf:
144164
case Builtin::BI__builtin_copysignf:
145-
case Builtin::BIcosf:
146-
case Builtin::BI__builtin_cosf:
147-
case Builtin::BIexpf:
148-
case Builtin::BI__builtin_expf:
149-
case Builtin::BIexp2f:
150-
case Builtin::BI__builtin_exp2f:
151165
case Builtin::BIfabsf:
152166
case Builtin::BI__builtin_fabsf:
153167
case Builtin::BIfloorf:
154168
case Builtin::BI__builtin_floorf:
155-
case Builtin::BIfmaf:
156-
case Builtin::BI__builtin_fmaf:
157169
case Builtin::BIfmaxf:
158170
case Builtin::BI__builtin_fmaxf:
159171
case Builtin::BIfminf:
160172
case Builtin::BI__builtin_fminf:
161-
case Builtin::BIfmodf:
162-
case Builtin::BI__builtin_fmodf:
163-
case Builtin::BIlogf:
164-
case Builtin::BI__builtin_logf:
165-
case Builtin::BIlog10f:
166-
case Builtin::BI__builtin_log10f:
167-
case Builtin::BIlog2f:
168-
case Builtin::BI__builtin_log2f:
169-
case Builtin::BIpowf:
170-
case Builtin::BI__builtin_powf:
173+
case Builtin::BInearbyintf:
174+
case Builtin::BI__builtin_nearbyintf:
171175
case Builtin::BIrintf:
172176
case Builtin::BI__builtin_rintf:
173177
case Builtin::BIroundf:
174178
case Builtin::BI__builtin_roundf:
175-
case Builtin::BIsinf:
176-
case Builtin::BI__builtin_sinf:
177-
case Builtin::BIsqrtf:
178-
case Builtin::BI__builtin_sqrtf:
179179
case Builtin::BItruncf:
180180
case Builtin::BI__builtin_truncf:
181181
case Builtin::BIlroundf:
182182
case Builtin::BI__builtin_lroundf:
183+
case Builtin::BI__builtin_fpclassify:
184+
case Builtin::BI__builtin_isfinite:
185+
case Builtin::BI__builtin_isinf:
186+
case Builtin::BI__builtin_isnormal:
183187
return false;
184188
default:
185189
break;

‎clang/test/SemaSYCL/unsupported_math.cpp

Lines changed: 33 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,14 @@
11
// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s
2-
2+
extern "C" float sinf(float);
3+
extern "C" float cosf(float);
4+
extern "C" float logf(float);
5+
extern "C" float ceilf(float);
6+
extern "C" float fabsf(float);
7+
extern "C" double sin(double);
8+
extern "C" double cos(double);
9+
extern "C" double log(double);
10+
extern "C" double ceil(double);
11+
extern "C" double fabs(double);
312
template <typename name, typename Func>
413
__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
514
kernelFunc();
@@ -9,9 +18,29 @@ int main() {
918
kernel<class kernel_function>([=]() {
1019
int acc[1] = {5};
1120
acc[0] *= 2;
12-
acc[0] += (int)__builtin_fabsf(-1.0f); // expected-error{{builtin is not supported on this target}}
13-
acc[0] += (int)__builtin_cosf(-1.0f); // expected-error{{builtin is not supported on this target}}
14-
acc[0] += (int)__builtin_powf(-1.0f, 10.0f); // expected-error{{builtin is not supported on this target}}
21+
acc[0] += (int)sinf(1.0f); // expected-no-error
22+
acc[0] += (int)sin(1.0); // expected-no-error
23+
acc[0] += (int)__builtin_sinf(1.0f); // expected-no-error
24+
acc[0] += (int)__builtin_sin(1.0); // expected-no-error
25+
acc[0] += (int)cosf(1.0f); // expected-no-error
26+
acc[0] += (int)cos(1.0); // expected-no-error
27+
acc[0] += (int)__builtin_cosf(1.0f); // expected-no-error
28+
acc[0] += (int)__builtin_cos(1.0); // expected-no-error
29+
acc[0] += (int)logf(1.0f); // expected-no-error
30+
acc[0] += (int)log(1.0); // expected-no-error
31+
acc[0] += (int)__builtin_logf(1.0f); // expected-no-error
32+
acc[0] += (int)__builtin_log(1.0); // expected-no-error
33+
acc[0] += (int)ceilf(1.0f); // expected-error{{builtin is not supported on this target}}
34+
acc[0] += (int)ceil(1.0); // expected-error{{builtin is not supported on this target}}
35+
acc[0] += (int)__builtin_ceilf(1.0f); // expected-error{{builtin is not supported on this target}}
36+
acc[0] += (int)__builtin_ceil(1.0); // expected-error{{builtin is not supported on this target}}
37+
acc[0] += (int)fabsf(-1.0f); // expected-error{{builtin is not supported on this target}}
38+
acc[0] += (int)fabs(-1.0); // expected-error{{builtin is not supported on this target}}
39+
acc[0] += (int)__builtin_fabsf(-1.0f); // expected-error{{builtin is not supported on this target}}
40+
acc[0] += (int)__builtin_fabs(-1.0); // expected-error{{builtin is not supported on this target}}
41+
acc[0] += (int)__builtin_fabsl(-1.0); // expected-error{{builtin is not supported on this target}}
42+
acc[0] += (int)__builtin_cosl(-1.0); // expected-error{{builtin is not supported on this target}}
43+
acc[0] += (int)__builtin_powl(-1.0, 10.0); // expected-error{{builtin is not supported on this target}}
1544
});
1645
return 0;
1746
}

‎sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

Lines changed: 116 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -23,10 +23,97 @@ or, in case of Windows:
2323
clang++ -fsycl main.obj %SYCL_INSTALL%/lib/libsycl-msvc.o -o a.exe
2424
2525
List of supported functions from C standard library:
26-
- assert macro (from <assert.h> or <cassert>)
27-
28-
NOTE: only the GNU glibc and Microsoft C libraries are currently
29-
supported.
26+
- assert macro (from <assert.h> or <cassert>)
27+
- logf, log (from <math.h> or <cmath>)
28+
- expf, exp (from <math.h> or <cmath>)
29+
- frexpf, frexp (from <math.h> or <cmath>)
30+
- ldexpf, ldexp (from <math.h> or <cmath>)
31+
- log10f, log10 (from <math.h> or <cmath>)
32+
- modff, modf (from <math.h> or <cmath>)
33+
- exp2f, exp2 (from <math.h> or <cmath>)
34+
- expm1f, expm1 (from <math.h> or <cmath>)
35+
- ilogbf, ilogb (from <math.h> or <cmath>)
36+
- log1pf, log1p (from <math.h> or <cmath>)
37+
- log2f, log2 (from <math.h> or <cmath>)
38+
- logbf, logb (from <math.h> or <cmath>)
39+
- sqrtf, sqrt (from <math.h> or <cmath>)
40+
- cbrtf, cbrt (from <math.h> or <cmath>)
41+
- hypotf, hypot (from <math.h> or <cmath>)
42+
- erff, erf (from <math.h> or <cmath>)
43+
- erfcf, erfc (from <math.h> or <cmath>)
44+
- tgammaf, tgamma (from <math.h> or <cmath>)
45+
- lgammaf, lgamma (from <math.h> or <cmath>)
46+
- fmodf, fmod (from <math.h> or <cmath>)
47+
- remainderf, remainder (from <math.h> or <cmath>)
48+
- remquof, remquo (from <math.h> or <cmath>)
49+
- nextafterf, nextafter (from <math.h> or <cmath>)
50+
- fdimf, fdim (from <math.h> or <cmath>)
51+
- fmaf, fma (from <math.h> or <cmath>)
52+
- sinf, sin (from <math.h> or <cmath>)
53+
- cosf, cos (from <math.h> or <cmath>)
54+
- tanf, tan (from <math.h> or <cmath>)
55+
- powf, pow (from <math.h> or <cmath>)
56+
- acosf, acos (from <math.h> or <cmath>)
57+
- asinf, asin (from <math.h> or <cmath>)
58+
- atanf, atan (from <math.h> or <cmath>)
59+
- atan2f, atan2 (from <math.h> or <cmath>)
60+
- coshf, cosh (from <math.h> or <cmath>)
61+
- sinhf, sinh (from <math.h> or <cmath>)
62+
- tanhf, tanh (from <math.h> or <cmath>)
63+
- acoshf, acosh (from <math.h> or <cmath>)
64+
- asinhf, asinh (from <math.h> or <cmath>)
65+
- atanhf, atanh (from <math.h> or <cmath>)
66+
- cimagf, cimag (from <complex.h>)
67+
- crealf, creal (from <complex.h>)
68+
- cargf, carg (from <complex.h>)
69+
- cabsf, cabs (from <complex.h>)
70+
- cprojf, cproj (from <complex.h>)
71+
- cexpf, cexp (from <complex.h>)
72+
- clogf, clog (from <complex.h>)
73+
- cpowf, cpow (from <complex.h>)
74+
- cpolarf, cpolar (from <complex.h>)
75+
- csqrtf, csqrt (from <complex.h>)
76+
- csinhf, csinh (from <complex.h>)
77+
- ccoshf, ccosh (from <complex.h>)
78+
- ctanhf, ctanh (from <complex.h>)
79+
- csinf, csin (from <complex.h>)
80+
- ccosf, ccos (from <complex.h>)
81+
- ctanf, ctan (from <complex.h>)
82+
- casinhf, casinh (from <complex.h>)
83+
- cacoshf, cacosh (from <complex.h>)
84+
- catanhf, catanh (from <complex.h>)
85+
- casinf, casin (from <complex.h>)
86+
- cacosf, cacos (from <complex.h>)
87+
- catanf, catan (from <complex.h>)
88+
89+
All functions are grouped into different device libraries based on
90+
functionalities. C and C++ standard library groups functions and
91+
classes by purpose(e.g. <math.h> for mathematical operations and
92+
transformations) and device library infrastructure uses this as
93+
a baseline.
94+
NOTE: Only the GNU glibc, Microsoft C libraries are currently
95+
supported. The device libraries for <math.h> and <complex.h> are
96+
ready for Linux and Windows support will be added in the future.
97+
Not all functions from <math.h> are supported right now, following
98+
math functions are not supported now:
99+
- abs
100+
- ceilf, ceil
101+
- copysignf, copysign
102+
- fabsf, fabs
103+
- floorf, floor
104+
- fmaxf, fmax
105+
- fminf, fmin
106+
- nextafterf, nextafter
107+
- rintf, rint
108+
- roundf, round
109+
- truncf, trunc
110+
- scalbnf, scalbn
111+
- nearbyintf, nearbyint
112+
- lrintf, lrint
113+
- nexttowardf, nexttoward
114+
- nanf, nan
115+
Device libraries can't support both single and double precision as some
116+
underlying device may not support double precision.
30117

31118
Example of usage
32119
================
@@ -58,6 +145,31 @@ Example of usage
58145
deviceQueue.wait_and_throw();
59146
}
60147
148+
149+
.. code: c++
150+
#include <math.h>
151+
#include <CL/sycl.hpp>
152+
153+
void device_sin_test() {
154+
cl::sycl::queue deviceQueue;
155+
cl::sycl::range<1> numOfItems{1};
156+
float result_f = -1.f;
157+
double result_d = -1.d;
158+
{
159+
cl::sycl::buffer<float, 1> buffer1(&result_f, numOfItems);
160+
cl::sycl::buffer<double, 1> buffer2(&result_d, numOfItems);
161+
deviceQueue.submit([&](cl::sycl::handler &cgh) {
162+
auto res_access1 = buffer1.get_access<sycl_write>(cgh);
163+
auto res_access2 = buffer2.get_access<sycl_write>(cgh);
164+
cgh.single_task<class DeviceSin>([=]() {
165+
res_access1[0] = sinf(0.f);
166+
res_access2[0] = sin(0.0);
167+
});
168+
});
169+
}
170+
assert((result_f == 0.f) && (result_d == 0.0));
171+
}
172+
61173
Frontend
62174
========
63175

0 commit comments

Comments
 (0)