Core: Fix SYCL/CUDA include-order compatibility libeigen/eigen!3406 Closes #2008
diff --git a/Eigen/src/Core/GenericPacketMath.h b/Eigen/src/Core/GenericPacketMath.h index 0be87b6..d05eb2b 100644 --- a/Eigen/src/Core/GenericPacketMath.h +++ b/Eigen/src/Core/GenericPacketMath.h
@@ -888,7 +888,7 @@ template <typename Scalar> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Scalar pscalar_from_float_bits(numext::uint32_t bits) { #if EIGEN_COMP_GNUC_STRICT && defined(__FINITE_MATH_ONLY__) && __FINITE_MATH_ONLY__ && \ - !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY) + !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) // GCC also needs the integer pattern hidden before bit_cast and widening. Use a memory operand so this works on // targets where EIGEN_OPTIMIZATION_BARRIER is intentionally unavailable. __asm__("" : "+m"(bits));
diff --git a/Eigen/src/Core/MathFunctions.h b/Eigen/src/Core/MathFunctions.h index eaa8920..73f65a1 100644 --- a/Eigen/src/Core/MathFunctions.h +++ b/Eigen/src/Core/MathFunctions.h
@@ -1279,7 +1279,7 @@ #endif #endif -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) #define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_char) \ @@ -1418,7 +1418,7 @@ return EIGEN_MATHFUNC_IMPL(hypot, Scalar)::run(x, y); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(hypot, hypot) #endif @@ -1427,7 +1427,7 @@ return EIGEN_MATHFUNC_IMPL(log1p, Scalar)::run(x); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log1p, log1p) #endif @@ -1449,7 +1449,7 @@ return internal::pow_impl<ScalarX, ScalarY>::run(x, y); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(pow, pow) #endif @@ -1466,7 +1466,7 @@ return internal::isfinite_impl(x); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isnan, isnan, bool) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isinf, isinf, bool) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isfinite, isfinite, bool) @@ -1497,7 +1497,7 @@ return internal::nearest_integer_impl<Scalar>::run_trunc(x); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(round, round) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(floor, floor) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(ceil, ceil) @@ -1598,7 +1598,7 @@ return x; } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sqrt, sqrt) #endif @@ -1626,7 +1626,7 @@ return internal::log_impl<T>::run(x); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log, log) #endif @@ -1663,7 +1663,7 @@ return x; } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(abs, abs) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(abs, fabs) #endif @@ -1748,7 +1748,7 @@ } #endif -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(exp, exp) #endif @@ -1794,7 +1794,7 @@ return internal::complex_exp2_impl<RealScalar>::run(x); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(exp2, exp2) #endif @@ -1818,7 +1818,7 @@ return static_cast<T>(ldexp(x, exponent)); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_float ldexp(const cl::sycl::cl_float& x, int exponent) { return cl::sycl::ldexp(x, exponent); @@ -1846,7 +1846,7 @@ return EIGEN_MATHFUNC_IMPL(expm1, Scalar)::run(x); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(expm1, expm1) #endif @@ -1868,7 +1868,7 @@ return cos(x); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cos, cos) #endif @@ -1890,7 +1890,7 @@ return sin(x); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sin, sin) #endif @@ -1912,7 +1912,7 @@ return tan(x); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tan, tan) #endif @@ -1940,7 +1940,7 @@ return static_cast<T>(acosh(x)); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acos, acos) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acosh, acosh) #endif @@ -1969,7 +1969,7 @@ return static_cast<T>(asinh(x)); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asin, asin) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asinh, asinh) #endif @@ -2004,7 +2004,7 @@ return static_cast<T>(atanh(x)); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atan, atan) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atanh, atanh) #endif @@ -2027,7 +2027,7 @@ return static_cast<T>(cosh(x)); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cosh, cosh) #endif @@ -2049,7 +2049,7 @@ return static_cast<T>(sinh(x)); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sinh, sinh) #endif @@ -2071,11 +2071,11 @@ return tanh(x); } -#if (!defined(EIGEN_GPUCC)) && EIGEN_FAST_MATH && !defined(SYCL_DEVICE_ONLY) +#if (!defined(EIGEN_GPUCC)) && EIGEN_FAST_MATH && !defined(EIGEN_SYCL_DEVICE_ONLY) EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float tanh(float x) { return internal::ptanh_float(x); } #endif -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tanh, tanh) #endif @@ -2097,7 +2097,7 @@ return fmod(a, b); } -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(fmod, fmod) #endif @@ -2154,7 +2154,7 @@ } #endif -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(nextafter, nextafter) #endif @@ -2170,7 +2170,7 @@ } #endif -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) #undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY #undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY #undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY
diff --git a/Eigen/src/Core/arch/Default/BFloat16.h b/Eigen/src/Core/arch/Default/BFloat16.h index 5fb7659..a04898a 100644 --- a/Eigen/src/Core/arch/Default/BFloat16.h +++ b/Eigen/src/Core/arch/Default/BFloat16.h
@@ -778,7 +778,7 @@ } EIGEN_DEVICE_FUNC inline bfloat16 fma(const bfloat16& a, const bfloat16& b, const bfloat16& c) { -#if defined(EIGEN_GPU_COMPILE_PHASE) || defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_GPU_COMPILE_PHASE) || defined(__SYCL_DEVICE_ONLY__) // Rounds twice; see half_impl::fma_round_to_odd. return bfloat16(numext::fma(static_cast<float>(a), static_cast<float>(b), static_cast<float>(c))); #else
diff --git a/Eigen/src/Core/arch/Default/Half.h b/Eigen/src/Core/arch/Default/Half.h index 6983b9c..e803565 100644 --- a/Eigen/src/Core/arch/Default/Half.h +++ b/Eigen/src/Core/arch/Default/Half.h
@@ -160,7 +160,7 @@ // CUDA GPU compile phase. -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(EIGEN_SYCL_DEVICE_ONLY) using __half_raw = cl::sycl::half; #endif @@ -755,7 +755,7 @@ return numext::bit_cast<numext::uint16_t>(h.x); #elif defined(EIGEN_HAS_BUILTIN_FLOAT16) return numext::bit_cast<numext::uint16_t>(h.x); -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(EIGEN_SYCL_DEVICE_ONLY) return numext::bit_cast<numext::uint16_t>(h); #else return h.x; @@ -1025,7 +1025,7 @@ #if defined(EIGEN_HAS_NATIVE_GPU_FP16) // Rounded once in hardware. return __hfma(::__half(a), ::__half(b), ::__half(c)); -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(__SYCL_DEVICE_ONLY__) // Rounds twice; see fma_round_to_odd. return half(numext::fma(static_cast<float>(a), static_cast<float>(b), static_cast<float>(c))); #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)
diff --git a/Eigen/src/Core/arch/SYCL/InteropHeaders.h b/Eigen/src/Core/arch/SYCL/InteropHeaders.h index 12c30d4..6b4907b 100644 --- a/Eigen/src/Core/arch/SYCL/InteropHeaders.h +++ b/Eigen/src/Core/arch/SYCL/InteropHeaders.h
@@ -78,7 +78,7 @@ }; }; -#ifdef SYCL_DEVICE_ONLY +#ifdef EIGEN_SYCL_DEVICE_ONLY #define SYCL_PACKET_TRAITS(packet_type, unpacket_type, lengths) \ template <> \ struct packet_traits<unpacket_type> : sycl_packet_traits<lengths> { \ @@ -129,7 +129,7 @@ template <typename PacketReturnType, int PacketSize> struct PacketWrapper; // This function should never get called on the device -#ifndef SYCL_DEVICE_ONLY +#ifndef EIGEN_SYCL_DEVICE_ONLY template <typename PacketReturnType, int PacketSize> struct PacketWrapper { typedef typename ::Eigen::internal::unpacket_traits<PacketReturnType>::type Scalar; @@ -147,7 +147,7 @@ } }; -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(EIGEN_SYCL_DEVICE_ONLY) template <typename PacketReturnType> struct PacketWrapper<PacketReturnType, 4> { typedef typename ::Eigen::internal::unpacket_traits<PacketReturnType>::type Scalar;
diff --git a/Eigen/src/Core/arch/SYCL/MathFunctions.h b/Eigen/src/Core/arch/SYCL/MathFunctions.h index 239aee9..72862c5 100644 --- a/Eigen/src/Core/arch/SYCL/MathFunctions.h +++ b/Eigen/src/Core/arch/SYCL/MathFunctions.h
@@ -32,7 +32,7 @@ // Make sure this is only available when targeting a GPU: we don't want to // introduce conflicts between these packet_traits definitions and the ones // we'll use on the host side (SSE, AVX, ...) -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) // Generic macro for unary SYCL math functions. #define SYCL_PACKET_FUNCTION(EIGEN_FUNC, SYCL_FUNC, PACKET) \
diff --git a/Eigen/src/Core/arch/SYCL/PacketMath.h b/Eigen/src/Core/arch/SYCL/PacketMath.h index 0c1bba8..753a91f 100644 --- a/Eigen/src/Core/arch/SYCL/PacketMath.h +++ b/Eigen/src/Core/arch/SYCL/PacketMath.h
@@ -30,7 +30,7 @@ namespace Eigen { namespace internal { -#ifdef SYCL_DEVICE_ONLY +#ifdef EIGEN_SYCL_DEVICE_ONLY #define SYCL_PLOAD(packet_type, AlignedType) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type pload##AlignedType<packet_type>( \ @@ -559,7 +559,7 @@ kernel.packet[1].x() = tmp; } -#endif // SYCL_DEVICE_ONLY +#endif // EIGEN_SYCL_DEVICE_ONLY } // end namespace internal
diff --git a/Eigen/src/Core/arch/SYCL/TypeCasting.h b/Eigen/src/Core/arch/SYCL/TypeCasting.h index 2c8fd51..1c3073f 100644 --- a/Eigen/src/Core/arch/SYCL/TypeCasting.h +++ b/Eigen/src/Core/arch/SYCL/TypeCasting.h
@@ -29,7 +29,7 @@ namespace Eigen { namespace internal { -#ifdef SYCL_DEVICE_ONLY +#ifdef EIGEN_SYCL_DEVICE_ONLY template <> struct type_casting_traits<float, int> { enum { VectorizedCast = 1, SrcCoeffRatio = 1, TgtCoeffRatio = 1 };
diff --git a/Eigen/src/Core/functors/UnaryFunctors.h b/Eigen/src/Core/functors/UnaryFunctors.h index c25cef0..22e126b 100644 --- a/Eigen/src/Core/functors/UnaryFunctors.h +++ b/Eigen/src/Core/functors/UnaryFunctors.h
@@ -1090,7 +1090,7 @@ template <typename Scalar, bool UseTypedPredicate = false> struct scalar_isnan_op { EIGEN_DEVICE_FUNC constexpr EIGEN_STRONG_INLINE bool operator()(const Scalar& a) const { -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) return numext::isnan(a); #else return numext::isnan EIGEN_NOT_A_MACRO(a); @@ -1101,7 +1101,7 @@ template <typename Scalar> struct scalar_isnan_op<Scalar, true> { EIGEN_DEVICE_FUNC constexpr EIGEN_STRONG_INLINE Scalar operator()(const Scalar& a) const { -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) return (numext::isnan(a) ? ptrue(a) : pzero(a)); #else return (numext::isnan EIGEN_NOT_A_MACRO(a) ? ptrue(a) : pzero(a)); @@ -1125,7 +1125,7 @@ template <typename Scalar, bool UseTypedPredicate = false> struct scalar_isinf_op { EIGEN_DEVICE_FUNC constexpr EIGEN_STRONG_INLINE bool operator()(const Scalar& a) const { -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) return numext::isinf(a); #else return (numext::isinf)(a); @@ -1136,7 +1136,7 @@ template <typename Scalar> struct scalar_isinf_op<Scalar, true> { EIGEN_DEVICE_FUNC constexpr EIGEN_STRONG_INLINE Scalar operator()(const Scalar& a) const { -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) return (numext::isinf(a) ? ptrue(a) : pzero(a)); #else return (numext::isinf EIGEN_NOT_A_MACRO(a) ? ptrue(a) : pzero(a)); @@ -1159,7 +1159,7 @@ template <typename Scalar, bool UseTypedPredicate = false> struct scalar_isfinite_op { EIGEN_DEVICE_FUNC constexpr EIGEN_STRONG_INLINE bool operator()(const Scalar& a) const { -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) return numext::isfinite(a); #else return (numext::isfinite)(a); @@ -1170,7 +1170,7 @@ template <typename Scalar> struct scalar_isfinite_op<Scalar, true> { EIGEN_DEVICE_FUNC constexpr EIGEN_STRONG_INLINE Scalar operator()(const Scalar& a) const { -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) return (numext::isfinite(a) ? ptrue(a) : pzero(a)); #else return (numext::isfinite EIGEN_NOT_A_MACRO(a) ? ptrue(a) : pzero(a));
diff --git a/Eigen/src/Core/util/Macros.h b/Eigen/src/Core/util/Macros.h index 9ba7a07..d71162e 100644 --- a/Eigen/src/Core/util/Macros.h +++ b/Eigen/src/Core/util/Macros.h
@@ -561,12 +561,14 @@ #error "NVCC as the target platform for HIPCC is currently not supported." #endif -#if defined(__CUDACC__) && !defined(EIGEN_NO_CUDA) && !defined(__SYCL_DEVICE_ONLY__) +// SYCL CUDA kernels still need CUDA scalar math unless Eigen's SYCL backend is selected. +#if defined(__CUDACC__) && !defined(EIGEN_NO_CUDA) && (!defined(__SYCL_DEVICE_ONLY__) || !defined(EIGEN_USE_SYCL)) // Means the compiler is either nvcc or clang with CUDA enabled #define EIGEN_CUDACC __CUDACC__ #endif -#if defined(__CUDA_ARCH__) && !defined(EIGEN_NO_CUDA) && !defined(__SYCL_DEVICE_ONLY__) +#if defined(__CUDA_ARCH__) && !defined(EIGEN_NO_CUDA) && \ + (!defined(__SYCL_DEVICE_ONLY__) || (defined(__CUDACC__) && !defined(EIGEN_USE_SYCL))) // Means we are generating code for the device #define EIGEN_CUDA_ARCH __CUDA_ARCH__ #endif @@ -721,10 +723,12 @@ #endif #if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__) -// EIGEN_USE_SYCL is a user-defined macro while __SYCL_DEVICE_ONLY__ is a compiler-defined macro. -// In most cases we want to check if both macros are defined which can be done using the define below. +// Eigen's SYCL dispatch flag. SYCL_DEVICE_ONLY is kept for users, but AdaptiveCpp's headers also define it. +#define EIGEN_SYCL_DEVICE_ONLY +#ifndef SYCL_DEVICE_ONLY #define SYCL_DEVICE_ONLY #endif +#endif // Under fast-math flags (-ffinite-math-only in particular), clang attaches the // `nofpclass(nan inf)` attribute to every function argument and return value of floating-point @@ -735,7 +739,7 @@ // store/load pair where it is materialized. It uses a memory operand so that it works for any // packet type, including aggregates, and expands to nothing outside affected builds. #if EIGEN_COMP_CLANG && defined(__FINITE_MATH_ONLY__) && __FINITE_MATH_ONLY__ && !defined(EIGEN_GPU_COMPILE_PHASE) && \ - !defined(SYCL_DEVICE_ONLY) + !defined(__SYCL_DEVICE_ONLY__) #define EIGEN_FAST_MATH_CONSTANT_BARRIER(X) __asm__("" : "+m"(X)) #else #define EIGEN_FAST_MATH_CONSTANT_BARRIER(X) @@ -936,7 +940,7 @@ // adds attributes to maximize inlining. This should only be used when really necessary: in particular, the // __attribute__((always_inline)) used on GCC is often unnecessary and can severely harm compile times. #ifndef EIGEN_ALWAYS_INLINE -#if EIGEN_COMP_GNUC && !defined(SYCL_DEVICE_ONLY) +#if EIGEN_COMP_GNUC && !defined(EIGEN_SYCL_DEVICE_ONLY) #define EIGEN_ALWAYS_INLINE __attribute__((always_inline)) inline #else #define EIGEN_ALWAYS_INLINE EIGEN_STRONG_INLINE @@ -949,7 +953,7 @@ // (those with auto parameters), so we leave it empty and rely on the // optimizer to inline small lambda bodies at /O2. #ifndef EIGEN_LAMBDA_ALWAYS_INLINE -#if EIGEN_COMP_GNUC && !defined(SYCL_DEVICE_ONLY) +#if EIGEN_COMP_GNUC && !defined(EIGEN_SYCL_DEVICE_ONLY) #define EIGEN_LAMBDA_ALWAYS_INLINE __attribute__((always_inline)) #else #define EIGEN_LAMBDA_ALWAYS_INLINE @@ -975,7 +979,7 @@ // GPU stuff // Disable some features when compiling with GPU compilers (SYCL/HIPCC) -#if defined(SYCL_DEVICE_ONLY) || defined(EIGEN_HIP_DEVICE_COMPILE) +#if defined(__SYCL_DEVICE_ONLY__) || defined(EIGEN_HIP_DEVICE_COMPILE) // Do not try asserts on device code #ifndef EIGEN_NO_DEBUG #define EIGEN_NO_DEBUG @@ -987,13 +991,13 @@ #endif // No exceptions on device. -#if defined(SYCL_DEVICE_ONLY) || defined(EIGEN_GPU_COMPILE_PHASE) +#if defined(__SYCL_DEVICE_ONLY__) || defined(EIGEN_GPU_COMPILE_PHASE) #ifdef EIGEN_EXCEPTIONS #undef EIGEN_EXCEPTIONS #endif #endif -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) && !defined(EIGEN_GPUCC) #ifndef EIGEN_DONT_VECTORIZE #define EIGEN_DONT_VECTORIZE #endif @@ -1365,7 +1369,7 @@ #elif EIGEN_ARCH_i386_OR_x86_64 // General, SSE. #define EIGEN_OPTIMIZATION_BARRIER(X) __asm__("" : "+g,x"(X)); -#elif !defined(SYCL_DEVICE_ONLY) +#elif !defined(__SYCL_DEVICE_ONLY__) // Other architectures: a memory operand needs no target-specific register class. #define EIGEN_OPTIMIZATION_BARRIER(X) __asm__("" : "+m"(X)); #else @@ -1408,12 +1412,12 @@ template <typename T> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void volatile_barrier(T& x) { EIGEN_UNUSED_VARIABLE(x); -#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY) && \ +#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) && \ (EIGEN_COMP_GNUC || (defined(__clang__) && defined(_MSC_VER))) __asm__ __volatile__("" : "+m"(x)); #endif } -#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY) && \ +#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) && \ (EIGEN_COMP_GNUC || (defined(__clang__) && defined(_MSC_VER))) #if EIGEN_ARCH_i386_OR_x86_64 && EIGEN_COMP_NVHPC // nvc++ rejects a long double asm operand under every constraint, so its bytes go through memory instead. @@ -1657,7 +1661,7 @@ #define EIGEN_FINAL final // Wrapping #pragma unroll in a macro since it is required for SYCL -#if defined(SYCL_DEVICE_ONLY) +#if defined(__SYCL_DEVICE_ONLY__) #if defined(_MSC_VER) #define EIGEN_UNROLL_LOOP __pragma(unroll) #else
diff --git a/contrib/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h b/contrib/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h index c8cfb06..ba67256 100644 --- a/contrib/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h +++ b/contrib/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h
@@ -67,10 +67,11 @@ template <> struct lgamma_impl<float> { EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE float run(float x) { -#if !defined(EIGEN_GPU_COMPILE_PHASE) && defined(EIGEN_HAS_LGAMMA_R) && !defined(__APPLE__) +#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) && defined(EIGEN_HAS_LGAMMA_R) && \ + !defined(__APPLE__) int dummy; return ::lgammaf_r(x, &dummy); -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(EIGEN_SYCL_DEVICE_ONLY) return cl::sycl::lgamma(x); #else return ::lgammaf(x); @@ -81,10 +82,11 @@ template <> struct lgamma_impl<double> { EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE double run(double x) { -#if !defined(EIGEN_GPU_COMPILE_PHASE) && defined(EIGEN_HAS_LGAMMA_R) && !defined(__APPLE__) +#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) && defined(EIGEN_HAS_LGAMMA_R) && \ + !defined(__APPLE__) int dummy; return ::lgamma_r(x, &dummy); -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(EIGEN_SYCL_DEVICE_ONLY) return cl::sycl::lgamma(x); #else return ::lgamma(x); @@ -474,7 +476,7 @@ template <> struct erfc_impl<float> { EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE float run(const float x) { -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) return cl::sycl::erfc(x); #else return generic_fast_erfc<float>::run(x); @@ -485,7 +487,7 @@ template <> struct erfc_impl<double> { EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE double run(const double x) { -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) return cl::sycl::erfc(x); #else return generic_fast_erfc<double>::run(x); @@ -594,7 +596,7 @@ template <> struct erf_impl<float> { EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE float run(const float x) { -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) return cl::sycl::erf(x); #else return generic_fast_erf<float>::run(x); @@ -605,7 +607,7 @@ template <> struct erf_impl<double> { EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE double run(const double x) { -#if defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_SYCL_DEVICE_ONLY) return cl::sycl::erf(x); #else return generic_fast_erf<double>::run(x);
diff --git a/contrib/Eigen/src/Tensor/TensorDeviceDefault.h b/contrib/Eigen/src/Tensor/TensorDeviceDefault.h index efb4bb4..af5d272 100644 --- a/contrib/Eigen/src/Tensor/TensorDeviceDefault.h +++ b/contrib/Eigen/src/Tensor/TensorDeviceDefault.h
@@ -64,7 +64,7 @@ } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE size_t firstLevelCacheSize() const { -#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY) +#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) // Running on the host CPU return l1CacheSize(); #elif defined(EIGEN_HIP_DEVICE_COMPILE) @@ -77,7 +77,7 @@ } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE size_t lastLevelCacheSize() const { -#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY) +#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) // Running single threaded on the host CPU return l3CacheSize(); #elif defined(EIGEN_HIP_DEVICE_COMPILE)
diff --git a/contrib/Eigen/src/Tensor/TensorIntDiv.h b/contrib/Eigen/src/Tensor/TensorIntDiv.h index 73c7adf..0d1a553 100644 --- a/contrib/Eigen/src/Tensor/TensorIntDiv.h +++ b/contrib/Eigen/src/Tensor/TensorIntDiv.h
@@ -23,7 +23,7 @@ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t<sizeof(T) == 4, int> count_leading_zeros(const T val) { #ifdef EIGEN_GPU_COMPILE_PHASE return __clz(val); -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(EIGEN_SYCL_DEVICE_ONLY) return cl::sycl::clz(val); #elif EIGEN_COMP_MSVC unsigned long index; @@ -39,7 +39,7 @@ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t<sizeof(T) == 8, int> count_leading_zeros(const T val) { #ifdef EIGEN_GPU_COMPILE_PHASE return __clzll(val); -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(EIGEN_SYCL_DEVICE_ONLY) return static_cast<int>(cl::sycl::clz(val)); #elif EIGEN_COMP_MSVC && EIGEN_ARCH_x86_64 unsigned long index; @@ -76,7 +76,7 @@ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE uint32_t muluh(const uint32_t a, const T b) { #if defined(EIGEN_GPU_COMPILE_PHASE) return __umulhi(a, b); -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(EIGEN_SYCL_DEVICE_ONLY) return cl::sycl::mul_hi(a, static_cast<uint32_t>(b)); #else return (static_cast<uint64_t>(a) * b) >> 32; @@ -87,7 +87,7 @@ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE uint64_t muluh(const uint64_t a, const T b) { #if defined(EIGEN_GPU_COMPILE_PHASE) return __umul64hi(a, b); -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(EIGEN_SYCL_DEVICE_ONLY) return cl::sycl::mul_hi(a, static_cast<uint64_t>(b)); #elif EIGEN_COMP_MSVC && (EIGEN_ARCH_x86_64 || EIGEN_ARCH_ARM64) return __umulh(a, static_cast<uint64_t>(b)); @@ -111,7 +111,7 @@ template <typename T> struct DividerHelper<64, T> { static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE uint64_t computeMultiplier(const int log_div, const T divider) { -#if EIGEN_HAS_BUILTIN_INT128 && !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY) +#if EIGEN_HAS_BUILTIN_INT128 && !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) return static_cast<uint64_t>((static_cast<__uint128_t>(1) << (64 + log_div)) / static_cast<__uint128_t>(divider) - (static_cast<__uint128_t>(1) << 64) + 1); #else @@ -194,7 +194,7 @@ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE int divide(const int32_t n) const { #ifdef EIGEN_GPU_COMPILE_PHASE return __umulhi(magic, n) >> shift; -#elif defined(SYCL_DEVICE_ONLY) +#elif defined(EIGEN_SYCL_DEVICE_ONLY) return cl::sycl::mul_hi(magic, static_cast<uint32_t>(n)) >> shift; #else uint64_t v = static_cast<uint64_t>(magic) * static_cast<uint64_t>(n);
diff --git a/contrib/Eigen/src/Tensor/TensorMacros.h b/contrib/Eigen/src/Tensor/TensorMacros.h index 1bdf913..a4e2824 100644 --- a/contrib/Eigen/src/Tensor/TensorMacros.h +++ b/contrib/Eigen/src/Tensor/TensorMacros.h
@@ -29,7 +29,7 @@ #define EIGEN_SFINAE_ENABLE_IF(__condition__) std::enable_if_t<(__condition__), int> = 0 // Define a macro to use a reference on the host but a value on the device -#if defined(SYCL_DEVICE_ONLY) +#if defined(__SYCL_DEVICE_ONLY__) #define EIGEN_DEVICE_REF #else #define EIGEN_DEVICE_REF &
diff --git a/contrib/Eigen/src/Tensor/TensorMorphing.h b/contrib/Eigen/src/Tensor/TensorMorphing.h index eb7a364..94d71f0 100644 --- a/contrib/Eigen/src/Tensor/TensorMorphing.h +++ b/contrib/Eigen/src/Tensor/TensorMorphing.h
@@ -1030,7 +1030,7 @@ } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index clamp(Index value, Index min, Index max) { -#ifndef SYCL_DEVICE_ONLY +#ifndef EIGEN_SYCL_DEVICE_ONLY return numext::maxi(min, numext::mini(max, value)); #else return cl::sycl::clamp(value, min, max);
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 3932508..5683a7a 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt
@@ -685,11 +685,14 @@ unset(EIGEN_ADD_TEST_FILENAME_EXTENSION) endif() +ei_add_test(sycl_macro_collision) + if(EIGEN_TEST_SYCL) set(EIGEN_SYCL ON) include(SyclConfigureTesting) ei_add_test(sycl_basic) + ei_add_test(sycl_include_order) set(EIGEN_SYCL OFF) endif()
diff --git a/test/fp_control.h b/test/fp_control.h index e0e9d5d..63f8d84 100644 --- a/test/fp_control.h +++ b/test/fp_control.h
@@ -11,7 +11,7 @@ #include <Eigen/Core> -#if defined(EIGEN_GPU_COMPILE_PHASE) || defined(SYCL_DEVICE_ONLY) +#if defined(EIGEN_GPU_COMPILE_PHASE) || defined(__SYCL_DEVICE_ONLY__) #define EIGEN_TEST_DEVICE_COMPILE 1 #else #define EIGEN_TEST_DEVICE_COMPILE 0
diff --git a/test/jacobi.cpp b/test/jacobi.cpp index 70868aa..2616120 100644 --- a/test/jacobi.cpp +++ b/test/jacobi.cpp
@@ -145,7 +145,7 @@ for (int mode = 0; mode < 4; ++mode) { JacobiRotation<Scalar> rotation, withoutR; Scalar r; -#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY) && EIGEN_ARCH_i386_OR_x86_64 && \ +#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) && EIGEN_ARCH_i386_OR_x86_64 && \ (defined(__SSE__) || defined(_M_X64) || (defined(_M_IX86_FP) && _M_IX86_FP >= 1)) { struct RestoreMxcsr {
diff --git a/test/numext.cpp b/test/numext.cpp index 91cecd1..1386fca 100644 --- a/test/numext.cpp +++ b/test/numext.cpp
@@ -986,7 +986,7 @@ std::numeric_limits<Scalar>::quiet_NaN()}; const int count = sizeof(values) / sizeof(values[0]); for (int mode = 0; mode < 4; ++mode) { -#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY) && EIGEN_ARCH_i386_OR_x86_64 && \ +#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) && EIGEN_ARCH_i386_OR_x86_64 && \ defined(_MM_SET_DENORMALS_ZERO_MODE) ScopedFlushToZero restore; const unsigned mask = _MM_FLUSH_ZERO_MASK | _MM_DENORMALS_ZERO_MASK; @@ -1070,7 +1070,7 @@ ScopedFlushToZero flush; if (!flush.isSupported()) return; check_binary_sign_no_flush_values<Scalar>(); -#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY) && EIGEN_ARCH_i386_OR_x86_64 && \ +#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) && EIGEN_ARCH_i386_OR_x86_64 && \ defined(_MM_SET_DENORMALS_ZERO_MODE) _MM_SET_DENORMALS_ZERO_MODE(_MM_DENORMALS_ZERO_ON); check_binary_sign_no_flush_values<Scalar>();
diff --git a/test/optimization_barrier.cpp b/test/optimization_barrier.cpp index 0bbc04e..d88c8b9 100644 --- a/test/optimization_barrier.cpp +++ b/test/optimization_barrier.cpp
@@ -83,7 +83,7 @@ VERIFY(!svptest_any(svptrue_b64(), svcmpne_n_f64(svptrue_b64(), sizeless, 1.25))); #endif -#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY) && \ +#if !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(__SYCL_DEVICE_ONLY__) && \ (EIGEN_COMP_GNUC || (defined(__clang__) && defined(_MSC_VER)) || (EIGEN_COMP_MSVC && defined(_M_FP_FAST))) VERIFY_IS_EQUAL(rounded_by_barrier(1.25f), 1.0f); VERIFY_IS_EQUAL(rounded_by_barrier(1.25), 1.0);
diff --git a/test/sycl_basic.cpp b/test/sycl_basic.cpp index ad9ba04..577bba3 100644 --- a/test/sycl_basic.cpp +++ b/test/sycl_basic.cpp
@@ -351,7 +351,7 @@ void test_predux_any_float() { auto operation = [](const float* in, int* out) { -#ifdef SYCL_DEVICE_ONLY +#ifdef EIGEN_SYCL_DEVICE_ONLY const cl::sycl::cl_int4 int_none(0); const cl::sycl::cl_int4 int_some(0, -1, 0, 0); const cl::sycl::cl_int4 int_all(-1); @@ -370,7 +370,7 @@ void test_predux_any_half() { auto operation = [](const float* in, int* out) { -#ifdef SYCL_DEVICE_ONLY +#ifdef EIGEN_SYCL_DEVICE_ONLY const cl::sycl::cl_short8 short_none(0); const cl::sycl::cl_short8 short_some(0, -1, 0, 0, 0, 0, 0, 0); const cl::sycl::cl_short8 short_all(-1); @@ -389,7 +389,7 @@ void test_predux_any_double() { auto operation = [](const float* in, int* out) { -#ifdef SYCL_DEVICE_ONLY +#ifdef EIGEN_SYCL_DEVICE_ONLY const cl::sycl::cl_long2 long_none(0); const cl::sycl::cl_long2 long_some(0, -1); const cl::sycl::cl_long2 long_all(-1);
diff --git a/test/sycl_include_order.cpp b/test/sycl_include_order.cpp new file mode 100644 index 0000000..34ef38f --- /dev/null +++ b/test/sycl_include_order.cpp
@@ -0,0 +1,105 @@ +// SPDX-FileCopyrightText: The Eigen Authors +// SPDX-License-Identifier: MPL-2.0 + +// EIGEN_SUFFIXES;1;2 + +// Core use inside a SYCL kernel does not require EIGEN_USE_SYCL. +#if defined(EIGEN_TEST_PART_2) +#include <Eigen/Core> +#endif +#include <CL/sycl.hpp> +#include <Eigen/Core> +#include <contrib/Eigen/SpecialFunctions> + +template <typename Scalar> +bool sycl_first_vector_addition() { + using Vector = Eigen::Matrix<Scalar, 4, 1>; + Vector a, b, result, expected; + a << 1, 2, 3, 4; + b << 5, 6, 7, 8; + expected << 6, 8, 10, 12; + cl::sycl::queue queue; + { + cl::sycl::buffer<Scalar, 1> input_a(a.data(), cl::sycl::range<1>(4)); + cl::sycl::buffer<Scalar, 1> input_b(b.data(), cl::sycl::range<1>(4)); + cl::sycl::buffer<Scalar, 1> output(result.data(), cl::sycl::range<1>(4)); + queue.submit([&](cl::sycl::handler& handler) { + auto x = input_a.template get_access<cl::sycl::access::mode::read>(handler); + auto y = input_b.template get_access<cl::sycl::access::mode::read>(handler); + auto z = output.template get_access<cl::sycl::access::mode::discard_write>(handler); + handler.single_task([=]() { + Eigen::Map<const Vector> a_map(&x[0]), b_map(&y[0]); + Eigen::Map<Vector> result_map(&z[0]); + result_map = a_map + b_map; + }); + }); + queue.wait_and_throw(); + } + return result == expected; +} + +template <typename Scalar> +bool sycl_scalar_math() { + using Vector = Eigen::Matrix<Scalar, 4, 1>; + Vector input, result_sin, result_lgamma; + input << 1, 2, 3, 4; + cl::sycl::queue queue; + { + cl::sycl::buffer<Scalar, 1> input_buffer(input.data(), cl::sycl::range<1>(4)); + cl::sycl::buffer<Scalar, 1> output_sin(result_sin.data(), cl::sycl::range<1>(4)); + cl::sycl::buffer<Scalar, 1> output_lgamma(result_lgamma.data(), cl::sycl::range<1>(4)); + queue.submit([&](cl::sycl::handler& handler) { + auto x = input_buffer.template get_access<cl::sycl::access::mode::read>(handler); + auto y = output_sin.template get_access<cl::sycl::access::mode::discard_write>(handler); + auto z = output_lgamma.template get_access<cl::sycl::access::mode::discard_write>(handler); + handler.single_task([=]() { + Eigen::Map<const Vector> input_map(&x[0]); + Eigen::Map<Vector> sin_map(&y[0]); + sin_map = input_map.array().sin().matrix(); + for (Eigen::Index i = 0; i < 4; ++i) z[i] = Eigen::numext::lgamma(input_map[i]); + }); + }); + queue.wait_and_throw(); + } + const Scalar tolerance = Scalar(32) * Eigen::NumTraits<Scalar>::epsilon(); + for (Eigen::Index i = 0; i < 4; ++i) { + const Scalar expected_sin = std::sin(input[i]); + const Scalar expected_lgamma = std::lgamma(input[i]); + if (!(std::abs(result_sin[i] - expected_sin) <= tolerance * (Scalar(1) + std::abs(expected_sin)))) return false; + if (!(std::abs(result_lgamma[i] - expected_lgamma) <= tolerance * (Scalar(1) + std::abs(expected_lgamma)))) + return false; + } + return true; +} + +template <typename Scalar> +bool sycl_low_precision_fma() { + float input[] = {1.5f, 2.f, 0.5f}, result = 0; + cl::sycl::queue queue; + { + cl::sycl::buffer<float, 1> inputs(input, cl::sycl::range<1>(3)); + cl::sycl::buffer<float, 1> output(&result, cl::sycl::range<1>(1)); + queue.submit([&](cl::sycl::handler& handler) { + auto x = inputs.template get_access<cl::sycl::access::mode::read>(handler); + auto y = output.template get_access<cl::sycl::access::mode::discard_write>(handler); + handler.single_task([=]() { y[0] = float(Eigen::numext::fma(Scalar(x[0]), Scalar(x[1]), Scalar(x[2]))); }); + }); + queue.wait_and_throw(); + } + const float expected = + float(Scalar(std::fma(float(Scalar(input[0])), float(Scalar(input[1])), float(Scalar(input[2]))))); + return result == expected; +} + +// The shared test driver explicitly enables CUDA when __CUDACC__ is defined; +// this test must leave backend selection to Eigen, as a user translation unit does. +int main() { + if (!sycl_first_vector_addition<float>() || !sycl_scalar_math<float>() || !sycl_low_precision_fma<Eigen::half>() || + !sycl_low_precision_fma<Eigen::bfloat16>()) + return 1; + cl::sycl::queue queue; + if (queue.get_device().has(cl::sycl::aspect::fp64) && + (!sycl_first_vector_addition<double>() || !sycl_scalar_math<double>())) + return 1; + return 0; +}
diff --git a/test/sycl_macro_collision.cpp b/test/sycl_macro_collision.cpp new file mode 100644 index 0000000..b435002 --- /dev/null +++ b/test/sycl_macro_collision.cpp
@@ -0,0 +1,68 @@ +// SPDX-FileCopyrightText: The Eigen Authors +// SPDX-License-Identifier: MPL-2.0 + +// EIGEN_SUFFIXES;1;2 + +#if defined(EIGEN_TEST_PART_2) +// Older DPC++ NVPTX passes define __CUDA_ARCH__ without CUDA language mode. +#define __SYCL_DEVICE_ONLY__ // NOLINT(bugprone-reserved-identifier): simulate the compiler's device pass. +#define __CUDA_ARCH__ 500 // NOLINT(bugprone-reserved-identifier): simulate older DPC++ NVPTX targets. +#ifndef EIGEN_NO_ASSERTION_CHECKING +#define EIGEN_NO_ASSERTION_CHECKING +#endif +#endif + +// AdaptiveCpp defines this flag even without Eigen's explicit SYCL integration. +#define SYCL_DEVICE_ONLY +#include "main.h" +#include "fp_control.h" +#if !defined(EIGEN_TEST_PART_2) +#include <contrib/Eigen/Tensor> +#endif + +#ifdef EIGEN_SYCL_DEVICE_ONLY +#error Eigen SYCL integration requires EIGEN_USE_SYCL +#endif +#ifndef SYCL_DEVICE_ONLY +#error Eigen must preserve externally defined macros +#endif + +#if defined(EIGEN_TEST_PART_2) +#if defined(EIGEN_CUDA_ARCH) || defined(EIGEN_GPU_COMPILE_PHASE) +#error SYCL device compilation alone must not select CUDA intrinsic paths +#endif +static_assert(sizeof(EIGEN_MAKESTRING(EIGEN_DEVICE_FUNC)) == 1, + "Implicit SYCL use must not enable Eigen's flattening policy"); +#endif + +EIGEN_DECLARE_TEST(sycl_macro_collision) { +#if !defined(EIGEN_TEST_PART_2) && EIGEN_ARCH_x86_64 + ScopedFlushToZero fp_mode(FlushToZeroMode::None); + VERIFY(fp_mode.isSupported()); +#endif +#if defined(EIGEN_TEST_PART_2) +#if EIGEN_COMP_GNUC + VERIFY_IS_EQUAL(std::string(EIGEN_MAKESTRING(EIGEN_ALWAYS_INLINE)), + std::string("__attribute__((always_inline)) inline")); + VERIFY_IS_EQUAL(std::string(EIGEN_MAKESTRING(EIGEN_LAMBDA_ALWAYS_INLINE)), + std::string("__attribute__((always_inline))")); +#endif + const half half_a(33.f / 32), half_b(65.f / 64), half_c((std::numeric_limits<half>::denorm_min)()); + VERIFY_IS_EQUAL(float(numext::fma(half_a, half_b, half_c)), + float(half(numext::fma(float(half_a), float(half_b), float(half_c))))); +#endif + Vector4d a, b; + a << 1, 2, 3, 4; + b << 5, 6, 7, 8; + Vector4d expected; + expected << 6, 8, 10, 12; + VERIFY_IS_EQUAL((a + b).eval(), expected); + VERIFY_IS_EQUAL(numext::mini(2, 3), 2); + VERIFY_IS_EQUAL(numext::maxi(2.0, 3.0), 3.0); +#if !defined(EIGEN_TEST_PART_2) + Tensor<float, 1> tensor(3), result(3); + tensor.setConstant(2.0f); + result = tensor + tensor; + for (Index i = 0; i < result.size(); ++i) VERIFY_IS_EQUAL(result(i), 4.0f); +#endif +}