Skip to content

Commit

Permalink
Use SYCL experimental extension for reduction properties (#2211)
Browse files Browse the repository at this point in the history
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.
  • Loading branch information
antonwolfy authored Dec 4, 2024
1 parent 0bac735 commit ed39ea7
Show file tree
Hide file tree
Showing 2 changed files with 20 additions and 3 deletions.
21 changes: 19 additions & 2 deletions dpnp/backend/kernels/dpnp_krnl_common.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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>
Expand Down Expand Up @@ -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]) *
Expand Down
2 changes: 1 addition & 1 deletion dpnp/tests/third_party/cupy/creation_tests/test_ranges.py
Original file line number Diff line number Diff line change
Expand Up @@ -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()
Expand Down

0 comments on commit ed39ea7

Please sign in to comment.