Merge branch 'master' into take_along_axis

Author: vtavana

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)
 * The parameter `axis` in `dpnp.take_along_axis` function has now a default value of `-1` [#2442](https://github.com/IntelPython/dpnp/pull/2442)
 
 ### Fixed

dpnp/backend/extensions/window/kaiser.cpp

@@ -32,22 +32,6 @@
 
 #include <sycl/sycl.hpp>
 
-/**
- * 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 <sycl/ext/intel/math.hpp> 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 <sycl/ext/intel/math.hpp>
-#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);

dpnp/backend/kernels/elementwise_functions/i0.hpp

@@ -35,16 +35,28 @@
 #define __SYCL_COMPILER_BESSEL_I0_SUPPORT 20241208L
 #endif
 
-// Include <sycl/ext/intel/math.hpp> 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 <sycl/ext/intel/math.hpp> 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 <sycl/ext/intel/math.hpp>
 #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 <typename argT, typename resT>
 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<resT, sycl::half>) {
             return static_cast<resT>(cyl_bessel_i0<float>(float(x)));
         }