diff --git a/intern/cycles/kernel/device/oneapi/compat.h b/intern/cycles/kernel/device/oneapi/compat.h index d8234ee1400..4794d1dcef5 100644 --- a/intern/cycles/kernel/device/oneapi/compat.h +++ b/intern/cycles/kernel/device/oneapi/compat.h @@ -196,9 +196,15 @@ ccl_always_inline float3 make_float3(float x) * include oneAPI headers, which transitively include math.h headers which will cause redefinitions * of the math defines because math.h also uses them and having them defined before math.h include * is actually UB. */ -/* Use fast math functions - get them from sycl::native namespace for native math function - * implementations */ -#define cosf(x) sycl::native::cos(((float)(x))) +/* sycl::native::cos precision is not sufficient and -ffast-math lets + * the current DPC++ compiler overload sycl::cos with it. + * We work around this issue by directly calling the spirv implementation which + * provides greater precision. */ +#if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__) +# define cosf(x) __spirv_ocl_cos(((float)(x))) +#else +# define cosf(x) sycl::cos(((float)(x))) +#endif #define sinf(x) sycl::native::sin(((float)(x))) #define powf(x, y) sycl::native::powr(((float)(x)), ((float)(y))) #define tanf(x) sycl::native::tan(((float)(x)))