Move tensor tests from dpctl

.github/workflows/check-onemath.yaml

@@ -74,7 +74,7 @@ jobs:
         os: [ubuntu-22.04] # windows-2022 - no DFT support for Windows in oneMKL
 
     runs-on: ${{ matrix.os }}
-    timeout-minutes: 60
+    timeout-minutes: 120
 
     defaults:
       run:
@@ -133,6 +133,14 @@ jobs:
         if: env.rerun-tests-on-failure != 'true'
         run: |
           python -m pytest -ra --pyargs dpnp.tests
+        env:
+          SKIP_TENSOR_TESTS: 1
+          SYCL_CACHE_PERSISTENT: 1
+
+      - name: Run tensor tests
+        if: env.rerun-tests-on-failure != 'true'
+        run: |
+          python -m pytest -ra --pyargs dpnp.tests.tensor
         env:
           SYCL_CACHE_PERSISTENT: 1
 
@@ -150,6 +158,24 @@ jobs:
             mamba activate ${{ env.test-env-name }}
 
             python -m pytest -ra --pyargs dpnp.tests
+        env:
+          SKIP_TENSOR_TESTS: 1
+          SYCL_CACHE_PERSISTENT: 1
+
+      - name: ReRun tensor tests on Linux
+        if: env.rerun-tests-on-failure == 'true'
+        id: run_tensor_tests
+        uses: nick-fields/retry@ce71cc2ab81d554ebbe88c79ab5975992d79ba08 # v3.0.2
+        with:
+          timeout_minutes: ${{ env.rerun-tests-timeout }}
+          max_attempts: ${{ env.rerun-tests-max-attempts }}
+          retry_on: any
+          command: |
+            . $CONDA/etc/profile.d/conda.sh
+            . $CONDA/etc/profile.d/mamba.sh
+            mamba activate ${{ env.test-env-name }}
+
+            python -m pytest -ra --pyargs dpnp.tests.tensor
         env:
           SYCL_CACHE_PERSISTENT: 1
 
@@ -239,6 +265,14 @@ jobs:
         if: env.rerun-tests-on-failure != 'true'
         run: |
           python -m pytest -ra --pyargs dpnp.tests
+        env:
+          SKIP_TENSOR_TESTS: 1
+          SYCL_CACHE_PERSISTENT: 1
+
+      - name: Run tensor tests
+        if: env.rerun-tests-on-failure != 'true'
+        run: |
+          python -m pytest -ra --pyargs dpnp.tests.tensor
         env:
           SYCL_CACHE_PERSISTENT: 1
 
@@ -256,5 +290,23 @@ jobs:
             mamba activate ${{ env.test-env-name }}
 
             python -m pytest -ra --pyargs dpnp.tests
+        env:
+          SKIP_TENSOR_TESTS: 1
+          SYCL_CACHE_PERSISTENT: 1
+
+      - name: ReRun tensor tests on Linux
+        if: env.rerun-tests-on-failure == 'true'
+        id: run_tensor_tests_branch
+        uses: nick-fields/retry@ce71cc2ab81d554ebbe88c79ab5975992d79ba08 # v3.0.2
+        with:
+          timeout_minutes: ${{ env.rerun-tests-timeout }}
+          max_attempts: ${{ env.rerun-tests-max-attempts }}
+          retry_on: any
+          command: |
+            . $CONDA/etc/profile.d/conda.sh
+            . $CONDA/etc/profile.d/mamba.sh
+            mamba activate ${{ env.test-env-name }}
+
+            python -m pytest -ra --pyargs dpnp.tests.tensor
         env:
           SYCL_CACHE_PERSISTENT: 1

.github/workflows/conda-package.yml

@@ -37,7 +37,7 @@ jobs:
       actions: write
 
     runs-on: ${{ matrix.os }}
-    timeout-minutes: 80
+    timeout-minutes: 90
 
     defaults:
       run:
@@ -220,6 +220,7 @@ jobs:
       - name: Run tests
         if: env.rerun-tests-on-failure != 'true'
         run: |
+          export SKIP_TENSOR_TESTS=1
           if [[ "${{ matrix.python }}" == "${{ env.python-ver-test-all-dtypes }}" ]]; then
             export DPNP_TEST_ALL_INT_TYPES=1
             python -m pytest -ra --pyargs ${{ env.package-name }}.tests
@@ -239,6 +240,7 @@ jobs:
             . $CONDA/etc/profile.d/conda.sh
             . $CONDA/etc/profile.d/mamba.sh
             mamba activate ${{ env.test-env-name }}
+            export SKIP_TENSOR_TESTS=1
 
             if [[ "${{ matrix.python }}" == "${{ env.python-ver-test-all-dtypes }}" ]]; then
               export DPNP_TEST_ALL_INT_TYPES=1
@@ -247,6 +249,26 @@ jobs:
               python -m pytest -n auto -ra --pyargs ${{ env.package-name }}.tests
             fi
 
+      - name: Run tensor tests
+        if: env.rerun-tests-on-failure != 'true'
+        run: |
+          python -m pytest -n auto -ra --pyargs dpnp.tests.tensor
+
+      - name: Run tensor tests
+        if: env.rerun-tests-on-failure == 'true'
+        id: run_tests_tensor_linux
+        uses: nick-fields/retry@ce71cc2ab81d554ebbe88c79ab5975992d79ba08 # v3.0.2
+        with:
+          timeout_minutes: ${{ env.rerun-tests-timeout }}
+          max_attempts: ${{ env.rerun-tests-max-attempts }}
+          retry_on: any
+          command: |
+            . $CONDA/etc/profile.d/conda.sh
+            . $CONDA/etc/profile.d/mamba.sh
+            mamba activate ${{ env.test-env-name }}
+
+            python -m pytest -n auto -ra --pyargs dpnp.tests.tensor
+
   test_windows:
     name: Test
 
@@ -382,6 +404,7 @@ jobs:
         if: env.rerun-tests-on-failure != 'true'
         shell: pwsh
         run: |
+          $env:SKIP_TENSOR_TESTS=1
           if (${{ matrix.python }} -eq ${{ env.python-ver-test-all-dtypes }}) {
             $env:DPNP_TEST_ALL_INT_TYPES=1
             python -m pytest -ra --pyargs ${{ env.package-name }}.tests
@@ -399,13 +422,32 @@ jobs:
           retry_on: any
           shell: pwsh
           command: |
+            $env:SKIP_TENSOR_TESTS=1
             if ( ${{ matrix.python }} -eq ${{ env.python-ver-test-all-dtypes }} ) {
               $env:DPNP_TEST_ALL_INT_TYPES=1
               python -m pytest -ra --pyargs ${{ env.package-name }}.tests
             } else {
               python -m pytest -n auto -ra --pyargs ${{ env.package-name }}.tests
             }
 
+      - name: Run tensor tests
+        if: env.rerun-tests-on-failure != 'true'
+        shell: pwsh
+        run: |
+          python -m pytest -n auto -ra --pyargs dpnp.tests.tensor
+
+      - name: Run tensor tests
+        if: env.rerun-tests-on-failure == 'true'
+        id: run_tests_tensor_win
+        uses: nick-fields/retry@ce71cc2ab81d554ebbe88c79ab5975992d79ba08 # v3.0.2
+        with:
+          timeout_minutes: ${{ env.rerun-tests-timeout }}
+          max_attempts: ${{ env.rerun-tests-max-attempts }}
+          retry_on: any
+          shell: pwsh
+          command: |
+            python -m pytest -n auto -ra --pyargs dpnp.tests.tensor
+
   upload:
     name: Upload
 

.github/workflows/generate_coverage.yaml

@@ -130,6 +130,7 @@ jobs:
             conda activate coverage
             [ -f /opt/intel/oneapi/setvars.sh ] && source /opt/intel/oneapi/setvars.sh
             git clean -fxd
+            export SKIP_TENSOR_TESTS=1
             python scripts/gen_coverage.py
 
       - name: Total number of coverage attempts

dpnp/tensor/CMakeLists.txt

@@ -29,6 +29,17 @@
 
 find_package(Python COMPONENTS Development.Module)
 
+# Tensor-specific flags
+
+# dpctl doesn't add -fsycl globally
+# only to pybind11 module sources via add_sycl_to_target()
+string(REPLACE "-fsycl " "" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
+
+# Use LLD linker (dpctl sets this at root level)
+if(UNIX)
+    add_link_options("-fuse-ld=lld")
+endif()
+
 # Remove global coverage flags for tensor
 # use link-time only approach like dpctl
 if(DPNP_GENERATE_COVERAGE)
@@ -55,6 +66,7 @@ if(
     endif()
 endif()
 
+# Match dpctl warning flags
 # Suppress unused parameter warnings
 add_compile_options(-Wno-unused-parameter)
 

dpnp/tensor/_usmarray.pyx

@@ -37,7 +37,7 @@ import numpy as np
 from dpctl._backend cimport DPCTLSyclUSMRef
 from dpctl._sycl_device_factory cimport _cached_default_device
 
-import dpnp.tensor
+import dpnp
 
 from ._data_types import bool as dpt_bool
 from ._device import Device

dpnp/tensor/libtensor/include/kernels/elementwise_functions/round.hpp

@@ -116,6 +116,7 @@ template <typename T>
 struct RoundOutputType
 {
     using value_type = typename std::disjunction<
+        td_ns::TypeMapResultEntry<T, bool, sycl::half>,
         td_ns::TypeMapResultEntry<T, std::uint8_t>,
         td_ns::TypeMapResultEntry<T, std::uint16_t>,
         td_ns::TypeMapResultEntry<T, std::uint32_t>,

dpnp/tensor/libtensor/include/utils/sycl_utils.hpp

@@ -283,9 +283,21 @@ T custom_inclusive_scan_over_group(GroupT &&wg,
         const bool in_range = (lane_id < n_aggregates);
         const bool in_bounds = in_range && (lane_id > 0 || large_wg);
 
-        T __scan_val = (in_bounds)
-                           ? local_mem_acc[(offset + lane_id) * max_sgSize - 1]
-                           : identity;
+        // Here is a bug where IGC incorrectly optimized the below code:
+        // T __scan_val = (in_bounds)
+        //                 ? local_mem_acc[(offset + lane_id) * max_sgSize - 1]
+        //                 : identity;
+        // That causes `__scan_val` is not initialized with `identity` value:
+        //   wgs = 256, max_sgSize = 16   =>   n_aggregates = 16
+        //   wi = 0:   in_range = 1, in_bounds = 0   =>   __scan_val = identity
+        // The w/s adds SYCL atomic fence, since the explicit memory fence
+        // prevents reordering/elimination, while it will add slight overhead.
+        T __scan_val = identity;
+        sycl::atomic_fence(sycl::memory_order::relaxed,
+                           sycl::memory_scope::work_item);
+        if (in_bounds) {
+            __scan_val = local_mem_acc[(offset + lane_id) * max_sgSize - 1];
+        }
         for (std::uint32_t step = 1; step < sgSize; step *= 2) {
             const bool advanced_lane = (lane_id >= step);
             const std::uint32_t src_lane_id =

dpnp/tests/config.py

@@ -4,6 +4,7 @@
 float16_types = bool(os.getenv("DPNP_TEST_FLOAT_16", 0))
 complex_types = bool(os.getenv("DPNP_TEST_COMPLEX_TYPES", 0))
 bool_types = bool(os.getenv("DPNP_TEST_BOOL_TYPES", 0))
+skip_tensor_tests = bool(int(os.getenv("SKIP_TENSOR_TESTS", 0)))
 
 
 infra_warnings_enable = bool(os.getenv("DPNP_INFRA_WARNINGS_ENABLE", 0))

dpnp/tests/conftest.py

@@ -97,6 +97,10 @@ def pytest_configure(config):
     # Equivalent to norecursedirs = tests_perf
     config.addinivalue_line("norecursedirs", "tests_perf")
 
+    # Equivalent to norecursedirs = tests/tensor (conditional)
+    if dtype_config.skip_tensor_tests:
+        config.addinivalue_line("norecursedirs", "tests/tensor")
+
     # Register pytest markers
     config.addinivalue_line(
         "markers", "slow: marks tests as slow (deselect with '-m \"not slow\"')"

dpnp/tests/tensor/__init__.py

@@ -0,0 +1,31 @@
+# *****************************************************************************
+# Copyright (c) 2026, Intel Corporation
+# All rights reserved.
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+# - Redistributions of source code must retain the above copyright notice,
+#   this list of conditions and the following disclaimer.
+# - Redistributions in binary form must reproduce the above copyright notice,
+#   this list of conditions and the following disclaimer in the documentation
+#   and/or other materials provided with the distribution.
+# - Neither the name of the copyright holder nor the names of its contributors
+#   may be used to endorse or promote products derived from this software
+#   without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+# THE POSSIBILITY OF SUCH DAMAGE.
+# *****************************************************************************
+
+__doc__ = r"""
+Test suite for tensor functionality migrated from dpctl.
+Running test suite requires Cython and a working compiler."""

dpnp/tests/tensor/conftest.py

@@ -0,0 +1,31 @@
+# *****************************************************************************
+# Copyright (c) 2026, Intel Corporation
+# All rights reserved.
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+# - Redistributions of source code must retain the above copyright notice,
+#   this list of conditions and the following disclaimer.
+# - Redistributions in binary form must reproduce the above copyright notice,
+#   this list of conditions and the following disclaimer in the documentation
+#   and/or other materials provided with the distribution.
+# - Neither the name of the copyright holder nor the names of its contributors
+#   may be used to endorse or promote products derived from this software
+#   without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+# THE POSSIBILITY OF SUCH DAMAGE.
+# *****************************************************************************
+
+"""Configures pytest to discover helper/ module"""
+
+from dpnp.tests.conftest import suppress_invalid_numpy_warnings

dpnp/tests/tensor/elementwise/__init__.py

@@ -0,0 +1,32 @@
+# *****************************************************************************
+# Copyright (c) 2026, Intel Corporation
+# All rights reserved.
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+# - Redistributions of source code must retain the above copyright notice,
+#   this list of conditions and the following disclaimer.
+# - Redistributions in binary form must reproduce the above copyright notice,
+#   this list of conditions and the following disclaimer in the documentation
+#   and/or other materials provided with the distribution.
+# - Neither the name of the copyright holder nor the names of its contributors
+#   may be used to endorse or promote products derived from this software
+#   without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+# THE POSSIBILITY OF SUCH DAMAGE.
+# *****************************************************************************
+
+"""
+Collection of test and utility files for testing elementwise operations
+over :class:`dpnp.tensor.usm_ndarray`.
+"""