Skip to content

Commit 3241425

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Use numeric_limits instead of NAN And INFINITY macros (#19852)
This commit replaces the uses of NAN and INFINITY in the SYCL complex headers with std::numeric_limits<float>::quiet_NaN() and std::numeric_limits<float>::infinity() respectively. This avoids issues where the definition of the macros could cause issues with differing constexpr'ness between platforms. Fixes #19114. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 64928c5 commit 3241425

File tree

4 files changed

+143
-112
lines changed

4 files changed

+143
-112
lines changed

‎sycl/include/sycl/ext/oneapi/experimental/bfloat16_math.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -44,7 +44,7 @@ constexpr int num_elements_v = sycl::detail::num_elements<T>::value;
4444

4545
/******************* isnan ********************/
4646

47-
// According to bfloat16 format, NAN value's exponent field is 0xFF and
47+
// According to bfloat16 format, NaN value's exponent field is 0xFF and
4848
// significand has non-zero bits.
4949
template <typename T>
5050
std::enable_if_t<std::is_same_v<T, bfloat16>, bool> isnan(T x) {

‎sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp

Lines changed: 21 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -187,7 +187,8 @@ __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
187187
proj(const complex<_Tp> &__c) {
188188
complex<_Tp> __r = __c;
189189
if (sycl::isinf(__c.real()) || sycl::isinf(__c.imag()))
190-
__r = complex<_Tp>(INFINITY, sycl::copysign(_Tp(0), __c.imag()));
190+
__r = complex<_Tp>(std::numeric_limits<float>::infinity(),
191+
sycl::copysign(_Tp(0), __c.imag()));
191192
return __r;
192193
}
193194

@@ -214,16 +215,18 @@ __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
214215
typename std::enable_if_t<is_genfloat<_Tp>::value, complex<_Tp>>
215216
polar(const _Tp &__rho, const _Tp &__theta = _Tp()) {
216217
if (sycl::isnan(__rho) || sycl::signbit(__rho))
217-
return complex<_Tp>(_Tp(NAN), _Tp(NAN));
218+
return complex<_Tp>(_Tp(std::numeric_limits<float>::quiet_NaN()),
219+
_Tp(std::numeric_limits<float>::quiet_NaN()));
218220
if (sycl::isnan(__theta)) {
219221
if (sycl::isinf(__rho))
220222
return complex<_Tp>(__rho, __theta);
221223
return complex<_Tp>(__theta, __theta);
222224
}
223225
if (sycl::isinf(__theta)) {
224226
if (sycl::isinf(__rho))
225-
return complex<_Tp>(__rho, _Tp(NAN));
226-
return complex<_Tp>(_Tp(NAN), _Tp(NAN));
227+
return complex<_Tp>(__rho, _Tp(std::numeric_limits<float>::quiet_NaN()));
228+
return complex<_Tp>(_Tp(std::numeric_limits<float>::quiet_NaN()),
229+
_Tp(std::numeric_limits<float>::quiet_NaN()));
227230
}
228231
_Tp __x = __rho * sycl::cos(__theta);
229232
if (sycl::isnan(__x))
@@ -259,7 +262,8 @@ __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
259262
typename std::enable_if_t<is_genfloat<_Tp>::value, complex<_Tp>>
260263
sqrt(const complex<_Tp> &__x) {
261264
if (sycl::isinf(__x.imag()))
262-
return complex<_Tp>(_Tp(INFINITY), __x.imag());
265+
return complex<_Tp>(_Tp(std::numeric_limits<float>::infinity()),
266+
__x.imag());
263267
if (sycl::isinf(__x.real())) {
264268
if (__x.real() > _Tp(0))
265269
return complex<_Tp>(__x.real(), sycl::isnan(__x.imag())
@@ -288,7 +292,7 @@ __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
288292
__i = _Tp(1);
289293
} else if (__i == 0 || !sycl::isfinite(__i)) {
290294
if (sycl::isinf(__i))
291-
__i = _Tp(NAN);
295+
__i = _Tp(std::numeric_limits<float>::quiet_NaN());
292296
return complex<_Tp>(__x.real(), __i);
293297
}
294298
}
@@ -436,8 +440,9 @@ __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
436440
sycl::copysign(__pi / _Tp(2), __x.imag()));
437441
}
438442
if (sycl::fabs(__x.real()) == _Tp(1) && __x.imag() == _Tp(0)) {
439-
return complex<_Tp>(sycl::copysign(_Tp(INFINITY), __x.real()),
440-
sycl::copysign(_Tp(0), __x.imag()));
443+
return complex<_Tp>(
444+
sycl::copysign(_Tp(std::numeric_limits<float>::infinity()), __x.real()),
445+
sycl::copysign(_Tp(0), __x.imag()));
441446
}
442447
complex<_Tp> __z = log((_Tp(1) + __x) / (_Tp(1) - __x)) / _Tp(2);
443448
return complex<_Tp>(sycl::copysign(__z.real(), __x.real()),
@@ -451,9 +456,11 @@ __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
451456
typename std::enable_if_t<is_genfloat<_Tp>::value, complex<_Tp>>
452457
sinh(const complex<_Tp> &__x) {
453458
if (sycl::isinf(__x.real()) && !sycl::isfinite(__x.imag()))
454-
return complex<_Tp>(__x.real(), _Tp(NAN));
459+
return complex<_Tp>(__x.real(),
460+
_Tp(std::numeric_limits<float>::quiet_NaN()));
455461
if (__x.real() == 0 && !sycl::isfinite(__x.imag()))
456-
return complex<_Tp>(__x.real(), _Tp(NAN));
462+
return complex<_Tp>(__x.real(),
463+
_Tp(std::numeric_limits<float>::quiet_NaN()));
457464
if (__x.imag() == 0 && !sycl::isfinite(__x.real()))
458465
return __x;
459466
return complex<_Tp>(sycl::sinh(__x.real()) * sycl::cos(__x.imag()),
@@ -467,9 +474,11 @@ __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY
467474
typename std::enable_if_t<is_genfloat<_Tp>::value, complex<_Tp>>
468475
cosh(const complex<_Tp> &__x) {
469476
if (sycl::isinf(__x.real()) && !sycl::isfinite(__x.imag()))
470-
return complex<_Tp>(sycl::fabs(__x.real()), _Tp(NAN));
477+
return complex<_Tp>(sycl::fabs(__x.real()),
478+
_Tp(std::numeric_limits<float>::quiet_NaN()));
471479
if (__x.real() == 0 && !sycl::isfinite(__x.imag()))
472-
return complex<_Tp>(_Tp(NAN), __x.real());
480+
return complex<_Tp>(_Tp(std::numeric_limits<float>::quiet_NaN()),
481+
__x.real());
473482
if (__x.real() == 0 && __x.imag() == 0)
474483
return complex<_Tp>(_Tp(1), __x.imag());
475484
if (__x.imag() == 0 && !sycl::isfinite(__x.real()))

‎sycl/include/sycl/stl_wrappers/__sycl_complex_impl.hpp

Lines changed: 60 additions & 38 deletions
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,8 @@
1212
// This header defines device-side overloads of <complex> functions.
1313

1414
#ifdef __SYCL_DEVICE_ONLY__
15-
#include <cmath> // For INFINITY.
15+
16+
#include <limits>
1617

1718
// The 'sycl_device_only' attribute enables device-side overloading.
1819
#define __SYCL_DEVICE __attribute__((sycl_device_only, always_inline))
@@ -82,8 +83,9 @@ float __complex__ __mulsc3(float __a, float __b, float __c, float __d) {
8283
__recalc = 1.0f;
8384
}
8485
if (__recalc) {
85-
z = __SYCL_CMPLXF((INFINITY * (__a * __c - __b * __d)),
86-
(INFINITY * (__a * __d + __b * __c)));
86+
z = __SYCL_CMPLXF(
87+
(std::numeric_limits<float>::infinity() * (__a * __c - __b * __d)),
88+
(std::numeric_limits<float>::infinity() * (__a * __d + __b * __c)));
8789
}
8890
}
8991
return z;
@@ -131,8 +133,9 @@ double __complex__ __muldc3(double __a, double __b, double __c, double __d) {
131133
__recalc = 1.0;
132134
}
133135
if (__recalc) {
134-
z = __SYCL_CMPLX((INFINITY * (__a * __c - __b * __d)),
135-
(INFINITY * (__a * __d + __b * __c)));
136+
z = __SYCL_CMPLX(
137+
(std::numeric_limits<float>::infinity() * (__a * __c - __b * __d)),
138+
(std::numeric_limits<float>::infinity() * (__a * __d + __b * __c)));
136139
}
137140
}
138141
return z;
@@ -161,15 +164,19 @@ float __complex__ __divsc3(float __a, float __b, float __c, float __d) {
161164
z = __SYCL_CMPLXF(z_real, z_imag);
162165
if (__spirv_IsNan(z_real) && __spirv_IsNan(z_imag)) {
163166
if ((__denom == 0.0f) && (!__spirv_IsNan(__a) || !__spirv_IsNan(__b))) {
164-
z_real = __spirv_ocl_copysign(INFINITY, __c) * __a;
165-
z_imag = __spirv_ocl_copysign(INFINITY, __c) * __b;
167+
z_real =
168+
__spirv_ocl_copysign(std::numeric_limits<float>::infinity(), __c) *
169+
__a;
170+
z_imag =
171+
__spirv_ocl_copysign(std::numeric_limits<float>::infinity(), __c) *
172+
__b;
166173
z = __SYCL_CMPLXF(z_real, z_imag);
167174
} else if ((__spirv_IsInf(__a) || __spirv_IsInf(__b)) &&
168175
__spirv_IsFinite(__c) && __spirv_IsFinite(__d)) {
169176
__a = __spirv_ocl_copysign(__spirv_IsInf(__a) ? 1.0f : 0.0f, __a);
170177
__b = __spirv_ocl_copysign(__spirv_IsInf(__b) ? 1.0f : 0.0f, __b);
171-
z_real = INFINITY * (__a * __c + __b * __d);
172-
z_imag = INFINITY * (__b * __c - __a * __d);
178+
z_real = std::numeric_limits<float>::infinity() * (__a * __c + __b * __d);
179+
z_imag = std::numeric_limits<float>::infinity() * (__b * __c - __a * __d);
173180
z = __SYCL_CMPLXF(z_real, z_imag);
174181
} else if (__spirv_IsInf(__logbw) && __logbw > 0.0f &&
175182
__spirv_IsFinite(__a) && __spirv_IsFinite(__b)) {
@@ -203,15 +210,19 @@ double __complex__ __divdc3(double __a, double __b, double __c, double __d) {
203210
z = __SYCL_CMPLX(z_real, z_imag);
204211
if (__spirv_IsNan(z_real) && __spirv_IsNan(z_imag)) {
205212
if ((__denom == 0.0) && (!__spirv_IsNan(__a) || !__spirv_IsNan(__b))) {
206-
z_real = __spirv_ocl_copysign((double)INFINITY, __c) * __a;
207-
z_imag = __spirv_ocl_copysign((double)INFINITY, __c) * __b;
213+
z_real =
214+
__spirv_ocl_copysign(std::numeric_limits<double>::infinity(), __c) *
215+
__a;
216+
z_imag =
217+
__spirv_ocl_copysign(std::numeric_limits<double>::infinity(), __c) *
218+
__b;
208219
z = __SYCL_CMPLX(z_real, z_imag);
209220
} else if ((__spirv_IsInf(__a) || __spirv_IsInf(__b)) &&
210221
__spirv_IsFinite(__c) && __spirv_IsFinite(__d)) {
211222
__a = __spirv_ocl_copysign(__spirv_IsInf(__a) ? 1.0 : 0.0, __a);
212223
__b = __spirv_ocl_copysign(__spirv_IsInf(__b) ? 1.0 : 0.0, __b);
213-
z_real = INFINITY * (__a * __c + __b * __d);
214-
z_imag = INFINITY * (__b * __c - __a * __d);
224+
z_real = std::numeric_limits<float>::infinity() * (__a * __c + __b * __d);
225+
z_imag = std::numeric_limits<float>::infinity() * (__b * __c - __a * __d);
215226
z = __SYCL_CMPLX(z_real, z_imag);
216227
} else if (__spirv_IsInf(__logbw) && __logbw > 0.0 &&
217228
__spirv_IsFinite(__a) && __spirv_IsFinite(__b)) {
@@ -247,14 +258,16 @@ __SYCL_DEVICE_C
247258
float __complex__ cprojf(float __complex__ z) {
248259
float __complex__ r = z;
249260
if (__spirv_IsInf(crealf(z)) || __spirv_IsInf(cimagf(z)))
250-
r = __SYCL_CMPLXF(INFINITY, __spirv_ocl_copysign(0.0f, cimagf(z)));
261+
r = __SYCL_CMPLXF(std::numeric_limits<float>::infinity(),
262+
__spirv_ocl_copysign(0.0f, cimagf(z)));
251263
return r;
252264
}
253265
__SYCL_DEVICE_C
254266
double __complex__ cproj(double __complex__ z) {
255267
double __complex__ r = z;
256268
if (__spirv_IsInf(creal(z)) || __spirv_IsInf(cimag(z)))
257-
r = __SYCL_CMPLX(INFINITY, __spirv_ocl_copysign(0.0, cimag(z)));
269+
r = __SYCL_CMPLX(std::numeric_limits<float>::infinity(),
270+
__spirv_ocl_copysign(0.0, cimag(z)));
258271
return r;
259272
}
260273

@@ -275,7 +288,7 @@ float __complex__ cexpf(float __complex__ z) {
275288
return z;
276289
} else if (z_imag == 0.f || !__spirv_IsFinite(z_imag)) {
277290
if (__spirv_IsInf(z_imag))
278-
return __SYCL_CMPLXF(z_real, NAN);
291+
return __SYCL_CMPLXF(z_real, std::numeric_limits<float>::quiet_NaN());
279292
}
280293
}
281294

@@ -293,16 +306,18 @@ double __complex__ cexp(double __complex__ z) {
293306
z_imag = 1.0;
294307
} else if (z_imag == 0.0 || !__spirv_IsFinite(z_imag)) {
295308
if (__spirv_IsInf(z_imag))
296-
z_imag = NAN;
309+
z_imag = std::numeric_limits<float>::quiet_NaN();
297310
return __SYCL_CMPLX(z_real, z_imag);
298311
}
299312
} else if (__spirv_IsNan(z_real)) {
300313
if (z_imag == 0.0)
301314
return z;
302-
return __SYCL_CMPLX(NAN, NAN);
315+
return __SYCL_CMPLX(std::numeric_limits<float>::quiet_NaN(),
316+
std::numeric_limits<float>::quiet_NaN());
303317
} else if (__spirv_IsFinite(z_real) &&
304318
(__spirv_IsNan(z_imag) || __spirv_IsInf(z_imag))) {
305-
return __SYCL_CMPLX(NAN, NAN);
319+
return __SYCL_CMPLX(std::numeric_limits<float>::quiet_NaN(),
320+
std::numeric_limits<float>::quiet_NaN());
306321
}
307322
double __e = __spirv_ocl_exp(z_real);
308323
double ret_real = __e * __spirv_ocl_cos(z_imag);
@@ -354,16 +369,18 @@ double __complex__ cpow(double __complex__ x, double __complex__ y) {
354369
__SYCL_DEVICE_C
355370
float __complex__ cpolarf(float rho, float theta) {
356371
if (__spirv_IsNan(rho) || __spirv_SignBitSet(rho))
357-
return __SYCL_CMPLXF(NAN, NAN);
372+
return __SYCL_CMPLXF(std::numeric_limits<float>::quiet_NaN(),
373+
std::numeric_limits<float>::quiet_NaN());
358374
if (__spirv_IsNan(theta)) {
359375
if (__spirv_IsInf(rho))
360376
return __SYCL_CMPLXF(rho, theta);
361377
return __SYCL_CMPLXF(theta, theta);
362378
}
363379
if (__spirv_IsInf(theta)) {
364380
if (__spirv_IsInf(rho))
365-
return __SYCL_CMPLXF(rho, NAN);
366-
return __SYCL_CMPLXF(NAN, NAN);
381+
return __SYCL_CMPLXF(rho, std::numeric_limits<float>::quiet_NaN());
382+
return __SYCL_CMPLXF(std::numeric_limits<float>::quiet_NaN(),
383+
std::numeric_limits<float>::quiet_NaN());
367384
}
368385
float x = rho * __spirv_ocl_cos(theta);
369386
if (__spirv_IsNan(x))
@@ -376,16 +393,18 @@ float __complex__ cpolarf(float rho, float theta) {
376393
__SYCL_DEVICE_C
377394
double __complex__ cpolar(double rho, double theta) {
378395
if (__spirv_IsNan(rho) || __spirv_SignBitSet(rho))
379-
return __SYCL_CMPLX(NAN, NAN);
396+
return __SYCL_CMPLX(std::numeric_limits<float>::quiet_NaN(),
397+
std::numeric_limits<float>::quiet_NaN());
380398
if (__spirv_IsNan(theta)) {
381399
if (__spirv_IsInf(rho))
382400
return __SYCL_CMPLX(rho, theta);
383401
return __SYCL_CMPLX(theta, theta);
384402
}
385403
if (__spirv_IsInf(theta)) {
386404
if (__spirv_IsInf(rho))
387-
return __SYCL_CMPLX(rho, NAN);
388-
return __SYCL_CMPLX(NAN, NAN);
405+
return __SYCL_CMPLX(rho, std::numeric_limits<float>::quiet_NaN());
406+
return __SYCL_CMPLX(std::numeric_limits<float>::quiet_NaN(),
407+
std::numeric_limits<float>::quiet_NaN());
389408
}
390409
double x = rho * __spirv_ocl_cos(theta);
391410
if (__spirv_IsNan(x))
@@ -401,7 +420,7 @@ float __complex__ csqrtf(float __complex__ z) {
401420
float z_real = crealf(z);
402421
float z_imag = cimagf(z);
403422
if (__spirv_IsInf(z_imag))
404-
return __SYCL_CMPLXF(INFINITY, z_imag);
423+
return __SYCL_CMPLXF(std::numeric_limits<float>::infinity(), z_imag);
405424
if (__spirv_IsInf(z_real)) {
406425
if (z_real > 0.0f)
407426
return __SYCL_CMPLXF(z_real, __spirv_IsNan(z_imag)
@@ -417,7 +436,7 @@ double __complex__ csqrt(double __complex__ z) {
417436
double z_real = creal(z);
418437
double z_imag = cimag(z);
419438
if (__spirv_IsInf(z_imag))
420-
return __SYCL_CMPLX(INFINITY, z_imag);
439+
return __SYCL_CMPLX(std::numeric_limits<float>::infinity(), z_imag);
421440
if (__spirv_IsInf(z_real)) {
422441
if (z_real > 0.0)
423442
return __SYCL_CMPLX(z_real, __spirv_IsNan(z_imag)
@@ -434,9 +453,9 @@ float __complex__ csinhf(float __complex__ z) {
434453
float z_real = crealf(z);
435454
float z_imag = cimagf(z);
436455
if (__spirv_IsInf(z_real) && !__spirv_IsFinite(z_imag))
437-
return __SYCL_CMPLXF(z_real, NAN);
456+
return __SYCL_CMPLXF(z_real, std::numeric_limits<float>::quiet_NaN());
438457
if (z_real == 0 && !__spirv_IsFinite(z_imag))
439-
return __SYCL_CMPLXF(z_real, NAN);
458+
return __SYCL_CMPLXF(z_real, std::numeric_limits<float>::quiet_NaN());
440459
if (z_imag == 0 && !__spirv_IsFinite(z_real))
441460
return z;
442461
return __SYCL_CMPLXF(__spirv_ocl_sinh(z_real) * __spirv_ocl_cos(z_imag),
@@ -447,9 +466,9 @@ double __complex__ csinh(double __complex__ z) {
447466
double z_real = creal(z);
448467
double z_imag = cimag(z);
449468
if (__spirv_IsInf(z_real) && !__spirv_IsFinite(z_imag))
450-
return __SYCL_CMPLX(z_real, NAN);
469+
return __SYCL_CMPLX(z_real, std::numeric_limits<float>::quiet_NaN());
451470
if (z_real == 0 && !__spirv_IsFinite(z_imag))
452-
return __SYCL_CMPLX(z_real, NAN);
471+
return __SYCL_CMPLX(z_real, std::numeric_limits<float>::quiet_NaN());
453472
if (z_imag == 0 && !__spirv_IsFinite(z_real))
454473
return z;
455474
return __SYCL_CMPLX(__spirv_ocl_sinh(z_real) * __spirv_ocl_cos(z_imag),
@@ -461,9 +480,10 @@ float __complex__ ccoshf(float __complex__ z) {
461480
float z_real = crealf(z);
462481
float z_imag = cimagf(z);
463482
if (__spirv_IsInf(z_real) && !__spirv_IsFinite(z_imag))
464-
return __SYCL_CMPLXF(__spirv_ocl_fabs(z_real), NAN);
483+
return __SYCL_CMPLXF(__spirv_ocl_fabs(z_real),
484+
std::numeric_limits<float>::quiet_NaN());
465485
if (z_real == 0 && !__spirv_IsFinite(z_imag))
466-
return __SYCL_CMPLXF(NAN, z_real);
486+
return __SYCL_CMPLXF(std::numeric_limits<float>::quiet_NaN(), z_real);
467487
if (z_real == 0 && z_imag == 0)
468488
return __SYCL_CMPLXF(1.0f, z_imag);
469489
if (z_imag == 0 && !__spirv_IsFinite(z_real))
@@ -476,9 +496,10 @@ double __complex__ ccosh(double __complex__ z) {
476496
double z_real = creal(z);
477497
double z_imag = cimag(z);
478498
if (__spirv_IsInf(z_real) && !__spirv_IsFinite(z_imag))
479-
return __SYCL_CMPLX(__spirv_ocl_fabs(z_real), NAN);
499+
return __SYCL_CMPLX(__spirv_ocl_fabs(z_real),
500+
std::numeric_limits<float>::quiet_NaN());
480501
if (z_real == 0 && !__spirv_IsFinite(z_imag))
481-
return __SYCL_CMPLX(NAN, z_real);
502+
return __SYCL_CMPLX(std::numeric_limits<float>::quiet_NaN(), z_real);
482503
if (z_real == 0 && z_imag == 0)
483504
return __SYCL_CMPLX(1.0f, z_imag);
484505
if (z_imag == 0 && !__spirv_IsFinite(z_real))
@@ -790,8 +811,9 @@ float __complex__ catanhf(float __complex__ z) {
790811
return __SYCL_CMPLXF(__spirv_ocl_copysign(0.0f, z_real),
791812
__spirv_ocl_copysign(__pi / 2.0f, z_imag));
792813
if (__spirv_ocl_fabs(z_real) == 1.0f && z_imag == 0.0f)
793-
return __SYCL_CMPLXF(__spirv_ocl_copysign(INFINITY, z_real),
794-
__spirv_ocl_copysign(0.0f, z_imag));
814+
return __SYCL_CMPLXF(
815+
__spirv_ocl_copysign(std::numeric_limits<float>::infinity(), z_real),
816+
__spirv_ocl_copysign(0.0f, z_imag));
795817
float __complex__ t1 = 1.0f + z;
796818
float __complex__ t2 = 1.0f - z;
797819
float __complex__ t3 =
@@ -820,7 +842,7 @@ double __complex__ catanh(double __complex__ z) {
820842
__spirv_ocl_copysign(__pi / 2.0, z_imag));
821843
if (__spirv_ocl_fabs(z_real) == 1.0 && z_imag == 0.0)
822844
return __SYCL_CMPLX(
823-
__spirv_ocl_copysign(static_cast<double>(INFINITY), z_real),
845+
__spirv_ocl_copysign(std::numeric_limits<double>::infinity(), z_real),
824846
__spirv_ocl_copysign(0.0, z_imag));
825847
double __complex__ t1 = 1.0 + z;
826848
double __complex__ t2 = 1.0 - z;

0 commit comments

Comments
 (0)