diff options
Diffstat (limited to 'intern/cycles/util/math.h')
| -rw-r--r-- | intern/cycles/util/math.h | 18 |
1 files changed, 15 insertions, 3 deletions
diff --git a/intern/cycles/util/math.h b/intern/cycles/util/math.h index f1f627588c5..af2f1ea092d 100644 --- a/intern/cycles/util/math.h +++ b/intern/cycles/util/math.h @@ -79,7 +79,7 @@ CCL_NAMESPACE_BEGIN /* Scalar */ -#ifndef __HIP__ +#if !defined(__HIP__) && !defined(__KERNEL_ONEAPI__) # ifdef _WIN32 ccl_device_inline float fmaxf(float a, float b) { @@ -92,12 +92,18 @@ ccl_device_inline float fminf(float a, float b) } # endif /* _WIN32 */ -#endif /* __HIP__ */ +#endif /* __HIP__, __KERNEL_ONEAPI__ */ -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) +# ifndef __KERNEL_ONEAPI__ using std::isfinite; using std::isnan; using std::sqrt; +# else +using sycl::sqrt; +# define isfinite(x) sycl::isfinite((x)) +# define isnan(x) sycl::isnan((x)) +# endif ccl_device_inline int abs(int x) { @@ -793,6 +799,8 @@ ccl_device_inline uint popcount(uint x) return i & 1; } # endif +#elif defined(__KERNEL_ONEAPI__) +# define popcount(x) sycl::popcount(x) #elif defined(__KERNEL_HIP__) /* Use popcll to support 64-bit wave for pre-RDNA AMD GPUs */ # define popcount(x) __popcll(x) @@ -806,6 +814,8 @@ ccl_device_inline uint count_leading_zeros(uint x) return __clz(x); #elif defined(__KERNEL_METAL__) return clz(x); +#elif defined(__KERNEL_ONEAPI__) + return sycl::clz(x); #else assert(x != 0); # ifdef _MSC_VER @@ -824,6 +834,8 @@ ccl_device_inline uint count_trailing_zeros(uint x) return (__ffs(x) - 1); #elif defined(__KERNEL_METAL__) return ctz(x); +#elif defined(__KERNEL_ONEAPI__) + return sycl::ctz(x); #else assert(x != 0); # ifdef _MSC_VER |