From ed39ea758124c3babfa983b2bceec1d9a64969ce Mon Sep 17 00:00:00 2001
From: Anton <100830759+antonwolfy@users.noreply.github.com>
Date: Wed, 4 Dec 2024 18:34:12 +0100
Subject: [PATCH] Use SYCL experimental extension for reduction properties
 (#2211)

Recent changes in DPC++ 2025.1 compiler broke DPNP compilation. It seems
the legacy DPC++ compiler behavior was affected by introducing SYCL
experimental extension for reduction properties.

Since the compilation issue occurred for a backend function which will
be removed once #2183 is merged, that PR proposes to temporary w/a the
issue and to reuse the extension for reduction properties, which seems
working.
---
 dpnp/backend/kernels/dpnp_krnl_common.cpp     | 21 +++++++++++++++++--
 .../cupy/creation_tests/test_ranges.py        |  2 +-
 2 files changed, 20 insertions(+), 3 deletions(-)

diff --git a/dpnp/backend/kernels/dpnp_krnl_common.cpp b/dpnp/backend/kernels/dpnp_krnl_common.cpp
index c7fb77ae4f5..0c25716a236 100644
--- a/dpnp/backend/kernels/dpnp_krnl_common.cpp
+++ b/dpnp/backend/kernels/dpnp_krnl_common.cpp
@@ -35,11 +35,23 @@
 #include "queue_sycl.hpp"
 #include <dpnp_iface.hpp>
 
+/**
+ * Version of SYCL DPC++ 2025.1 compiler where support of
+ * sycl::ext::oneapi::experimental::properties was added.
+ */
+#ifndef __SYCL_COMPILER_REDUCTION_PROPERTIES_SUPPORT
+#define __SYCL_COMPILER_REDUCTION_PROPERTIES_SUPPORT 20241129
+#endif
+
 namespace mkl_blas = oneapi::mkl::blas;
 namespace mkl_blas_cm = oneapi::mkl::blas::column_major;
 namespace mkl_blas_rm = oneapi::mkl::blas::row_major;
 namespace mkl_lapack = oneapi::mkl::lapack;
 
+#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_REDUCTION_PROPERTIES_SUPPORT
+namespace syclex = sycl::ext::oneapi::experimental;
+#endif
+
 template <typename _KernelNameSpecialization1,
           typename _KernelNameSpecialization2,
           typename _KernelNameSpecialization3>
@@ -78,8 +90,13 @@ sycl::event dot(sycl::queue &queue,
             cgh.parallel_for(
                 sycl::range<1>{size},
                 sycl::reduction(
-                    result_out, std::plus<_DataType_output>(),
-                    sycl::property::reduction::initialize_to_identity{}),
+                    result_out, sycl::plus<_DataType_output>(),
+#if __SYCL_COMPILER_VERSION >= __SYCL_COMPILER_REDUCTION_PROPERTIES_SUPPORT
+                    syclex::properties(syclex::initialize_to_identity)
+#else
+                    sycl::property::reduction::initialize_to_identity {}
+#endif
+                        ),
                 [=](sycl::id<1> idx, auto &sum) {
                     sum += static_cast<_DataType_output>(
                                input1_in[idx * input1_strides]) *
diff --git a/dpnp/tests/third_party/cupy/creation_tests/test_ranges.py b/dpnp/tests/third_party/cupy/creation_tests/test_ranges.py
index bb93179ced7..72249622334 100644
--- a/dpnp/tests/third_party/cupy/creation_tests/test_ranges.py
+++ b/dpnp/tests/third_party/cupy/creation_tests/test_ranges.py
@@ -172,7 +172,7 @@ def test_linspace_float_overflow(self, xp):
         dtype = cupy.default_float_type()
         return xp.linspace(0.0, xp.finfo(dtype).max / 5, 10, dtype=dtype)
 
-    @testing.numpy_cupy_allclose()
+    @testing.numpy_cupy_allclose(rtol={numpy.float32: 1e-6, "default": 1e-7})
     def test_linspace_float_underflow(self, xp):
         # find minimum subnormal number
         dtype = cupy.default_float_type()