Merge branch 'master' into impl_of_interp

Author: antonwolfy

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
 

conda-recipe/meta.yaml

@@ -13,7 +13,7 @@ requirements:
       - python
       - setuptools
       - numpy >=1.23
-      - cython <3.1
+      - cython
       - cmake >=3.21
       - ninja
       - git

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)));
         }

dpnp/dpnp_algo/dpnp_arraycreation.py

@@ -32,8 +32,8 @@
 import numpy
 
 import dpnp
-import dpnp.dpnp_utils as utils
 from dpnp.dpnp_array import dpnp_array
+from dpnp.dpnp_utils import get_usm_allocations, map_dtype_to_device
 
 __all__ = [
     "dpnp_geomspace",
@@ -60,7 +60,7 @@ def dpnp_geomspace(
     endpoint=True,
     axis=0,
 ):
-    usm_type_alloc, sycl_queue_alloc = utils.get_usm_allocations([start, stop])
+    usm_type_alloc, sycl_queue_alloc = get_usm_allocations([start, stop])
 
     if sycl_queue is None and device is None:
         sycl_queue = sycl_queue_alloc
@@ -77,7 +77,7 @@ def dpnp_geomspace(
     stop = _as_usm_ndarray(stop, _usm_type, sycl_queue_normalized)
 
     dt = numpy.result_type(start, stop, float(num))
-    dt = utils.map_dtype_to_device(dt, sycl_queue_normalized.sycl_device)
+    dt = map_dtype_to_device(dt, sycl_queue_normalized.sycl_device)
     if dtype is None:
         dtype = dt
 
@@ -144,7 +144,7 @@ def dpnp_linspace(
     retstep=False,
     axis=0,
 ):
-    usm_type_alloc, sycl_queue_alloc = utils.get_usm_allocations([start, stop])
+    usm_type_alloc, sycl_queue_alloc = get_usm_allocations([start, stop])
 
     if sycl_queue is None and device is None:
         sycl_queue = sycl_queue_alloc
@@ -164,7 +164,7 @@ def dpnp_linspace(
         stop = _as_usm_ndarray(stop, _usm_type, sycl_queue_normalized)
 
     dt = numpy.result_type(start, stop, float(num))
-    dt = utils.map_dtype_to_device(dt, sycl_queue_normalized.sycl_device)
+    dt = map_dtype_to_device(dt, sycl_queue_normalized.sycl_device)
     if dtype is None:
         dtype = dt
 
@@ -251,7 +251,7 @@ def dpnp_logspace(
     axis=0,
 ):
     if not dpnp.isscalar(base):
-        usm_type_alloc, sycl_queue_alloc = utils.get_usm_allocations(
+        usm_type_alloc, sycl_queue_alloc = get_usm_allocations(
             [start, stop, base]
         )
 

dpnp/dpnp_utils/dpnp_algo_utils.pyx

@@ -449,6 +449,7 @@ cdef tuple get_common_usm_allocation(dpnp_descriptor x1, dpnp_descriptor x2):
     return (common_sycl_queue.sycl_device, common_usm_type, common_sycl_queue)
 
 
+@cython.linetrace(False)
 cdef (DPNPFuncType, void *) get_ret_type_and_func(DPNPFuncData kernel_data,
                                                   cpp_bool has_aspect_fp64):
     """

environments/build_with_oneapi.yml

@@ -3,7 +3,7 @@ channels:
   - conda-forge
 dependencies:
   - cmake
-  - cython <3.1
+  - cython
   - ninja
   - numpy
   - pytest