diff --git a/CHANGELOG.md b/CHANGELOG.md index c501ef7e5003..033da8fd0822 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -28,6 +28,7 @@ This release achieves 100% compliance with Python Array API specification (revis * Updated Python Array API specification version supported to `2024.12` [#2416](https://github.com/IntelPython/dpnp/pull/2416) * Removed `einsum_call` keyword from `dpnp.einsum_path` signature [#2421](https://github.com/IntelPython/dpnp/pull/2421) * Changed `"max dimensions"` to `None` in array API capabilities [#2432](https://github.com/IntelPython/dpnp/pull/2432) +* Updated kernel header `i0.hpp` to expose `cyl_bessel_i0` function depending on build target [#2440](https://github.com/IntelPython/dpnp/pull/2440) ### Fixed diff --git a/dpnp/backend/extensions/window/kaiser.cpp b/dpnp/backend/extensions/window/kaiser.cpp index fd86d3dda773..68ac1774c63e 100644 --- a/dpnp/backend/extensions/window/kaiser.cpp +++ b/dpnp/backend/extensions/window/kaiser.cpp @@ -32,22 +32,6 @@ #include -/** - * Version of SYCL DPC++ 2025.1 compiler where an issue with - * sycl::ext::intel::math::cyl_bessel_i0(x) is fully resolved. - */ -#ifndef __SYCL_COMPILER_BESSEL_I0_SUPPORT -#define __SYCL_COMPILER_BESSEL_I0_SUPPORT 20241208L -#endif - -// Include only when targeting Intel devices. -// This header relies on intel-specific types like _iml_half_internal, -// which are not supported on non-intel backends (e.g., CUDA, AMD) -#if defined(__SPIR__) && defined(__INTEL_LLVM_COMPILER) && \ - (__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_BESSEL_I0_SUPPORT) -#include -#endif - #include "../kernels/elementwise_functions/i0.hpp" namespace dpnp::extensions::window @@ -78,12 +62,7 @@ class KaiserFunctor void operator()(sycl::id<1> id) const { -#if defined(__SPIR__) && defined(__INTEL_LLVM_COMPILER) && \ - (__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_BESSEL_I0_SUPPORT) - using sycl::ext::intel::math::cyl_bessel_i0; -#else - using dpnp::kernels::i0::impl::cyl_bessel_i0; -#endif + using dpnp::kernels::i0::cyl_bessel_i0; const auto i = id.get(0); const T alpha = (N - 1) / T(2); diff --git a/dpnp/backend/kernels/elementwise_functions/i0.hpp b/dpnp/backend/kernels/elementwise_functions/i0.hpp index e474bcb41f00..995263b28e66 100644 --- a/dpnp/backend/kernels/elementwise_functions/i0.hpp +++ b/dpnp/backend/kernels/elementwise_functions/i0.hpp @@ -35,16 +35,28 @@ #define __SYCL_COMPILER_BESSEL_I0_SUPPORT 20241208L #endif -// Include only when targeting Intel devices. -// This header relies on intel-specific types like _iml_half_internal, -// which are not supported on non-intel backends (e.g., CUDA, AMD) -#if defined(__SPIR__) && defined(__INTEL_LLVM_COMPILER) && \ +/** + * Include only when targeting to Intel devices. + * This header relies on intel-specific types like _iml_half_internal, + * which are not suppose to work with other targets (e.g., CUDA, AMD). + */ +#if defined(__SPIR__) && defined(__INTEL_LLVM_COMPILER) +#define __SYCL_EXT_INTEL_MATH_SUPPORT +#endif + +#if defined(__SYCL_EXT_INTEL_MATH_SUPPORT) && \ (__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_BESSEL_I0_SUPPORT) #include #endif namespace dpnp::kernels::i0 { +#if defined(__SYCL_EXT_INTEL_MATH_SUPPORT) && \ + (__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_BESSEL_I0_SUPPORT) +using sycl::ext::intel::math::cyl_bessel_i0; + +#else + /** * The below implementation of Bessel function of order 0 * is based on the source code from https://github.com/gcc-mirror/gcc @@ -243,6 +255,10 @@ inline Tp cyl_bessel_i0(Tp x) } } // namespace impl +using impl::cyl_bessel_i0; + +#endif + template struct I0Functor { @@ -257,13 +273,6 @@ struct I0Functor resT operator()(const argT &x) const { -#if defined(__SPIR__) && defined(__INTEL_LLVM_COMPILER) && \ - (__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_BESSEL_I0_SUPPORT) - using sycl::ext::intel::math::cyl_bessel_i0; -#else - using impl::cyl_bessel_i0; -#endif - if constexpr (std::is_same_v) { return static_cast(cyl_bessel_i0(float(x))); }