diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml index 2fda6913bc5c..8ae37441bba2 100644 --- a/.github/workflows/conda-package.yml +++ b/.github/workflows/conda-package.yml @@ -15,6 +15,7 @@ env: test_arraycreation.py test_dot.py test_dparray.py + test_copy.py test_fft.py test_linalg.py test_logic.py @@ -23,12 +24,16 @@ env: test_random_state.py test_sort.py test_special.py + test_sycl_queue.py test_umath.py test_usm_type.py third_party/cupy/linalg_tests/test_product.py + third_party/cupy/logic_tests/test_comparison.py third_party/cupy/logic_tests/test_truth.py third_party/cupy/manipulation_tests/test_basic.py third_party/cupy/manipulation_tests/test_join.py + third_party/cupy/manipulation_tests/test_rearrange.py + third_party/cupy/manipulation_tests/test_transpose.py third_party/cupy/math_tests/test_explog.py third_party/cupy/math_tests/test_misc.py third_party/cupy/math_tests/test_trigonometric.py @@ -43,7 +48,7 @@ jobs: strategy: matrix: - python: ['3.8', '3.9', '3.10', '3.11'] + python: ['3.9', '3.10', '3.11'] os: [ubuntu-20.04, windows-latest] runs-on: ${{ matrix.os }} @@ -120,7 +125,7 @@ jobs: strategy: matrix: - python: ['3.8', '3.9', '3.10', '3.11'] + python: ['3.9', '3.10', '3.11'] os: [ubuntu-20.04, ubuntu-latest] continue-on-error: true @@ -221,7 +226,7 @@ jobs: strategy: matrix: - python: ['3.8', '3.9', '3.10', '3.11'] + python: ['3.9', '3.10', '3.11'] continue-on-error: true @@ -353,7 +358,7 @@ jobs: strategy: matrix: - python: ['3.8', '3.9', '3.10', '3.11'] + python: ['3.9', '3.10', '3.11'] os: [ubuntu-20.04, windows-latest] runs-on: ${{ matrix.os }} diff --git a/.github/workflows/generate_coverage.yaml b/.github/workflows/generate_coverage.yaml index 413835d336cf..776e521f66e0 100644 --- a/.github/workflows/generate_coverage.yaml +++ b/.github/workflows/generate_coverage.yaml @@ -49,8 +49,7 @@ jobs: conda list - name: Build dpnp with coverage run: | - python scripts/gen_coverage.py --pytest-opts="--ignore tests/test_random.py \ - --ignore tests/test_strides.py" + python scripts/gen_coverage.py --pytest-opts="--ignore tests/test_random.py" - name: Install coverall dependencies run: | sudo gem install coveralls-lcov diff --git a/CHANGELOG.md b/CHANGELOG.md index 01911ea4911c..a15f81807cc2 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -4,14 +4,66 @@ All notable changes to this project will be documented in this file. The format is based on [Keep a Changelog](https://keepachangelog.com/en/1.0.0/), and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html). -## [0.12.2] - TBA +## [0.13.0] - TBA ### Added +* Added implementation of flipping functions: `dpnp.flip`, `dpnp.fliplr` and `dpnp.flipud` [#1543](https://github.com/IntelPython/dpnp/pull/1543) +* Added implementation of `dpnp.rint` function through `dpnp.round` call [#1537](https://github.com/IntelPython/dpnp/pull/1537) +* Added in-place support for arithmetic operators [#1530](https://github.com/IntelPython/dpnp/pull/1530) +* Dropped build and uploading the package with `python=3.8` to `dppy/label/dev` channel of Anaconda [#1534](https://github.com/IntelPython/dpnp/pull/1534) +* Implemented build and uploading the package with `python=3.11` to `dppy/label/dev` channel of Anaconda [#1501](https://github.com/IntelPython/dpnp/pull/1501) +* Added the versioneer to compute a product version number [#1497](https://github.com/IntelPython/dpnp/pull/1497) +* Added `cython` support of `3.0.0` or above version [#1495](https://github.com/IntelPython/dpnp/pull/1495) + ### Changed +* Updated `Build from source` section in `README.md` to state all the required prerequisite packages [#1553](https://github.com/IntelPython/dpnp/pull/1553) +* Reworked `dpnp.hstack` and `dpnp.atleast_1d` through existing functions to get rid of falling back on NumPy [#1544](https://github.com/IntelPython/dpnp/pull/1544) +* Reworked `dpnp.asfarray` through existing functions to get rid of falling back on NumPy [#1542](https://github.com/IntelPython/dpnp/pull/1542) +* Converted from `raw` to `multi_ptr` with `address_space_cast` to adopt towards changes introduced in `SYCL 2020` [#1538](https://github.com/IntelPython/dpnp/pull/1538) +* Updated install instruction via `pip` [#1531](https://github.com/IntelPython/dpnp/pull/1531) +* Reworked `dpnp.copyto` through existing functions instead of a separate kernel [#1516](https://github.com/IntelPython/dpnp/pull/1516) +* Aligned default order value with NumPy in asarray-like functions [#1526](https://github.com/IntelPython/dpnp/pull/1526) +* Created unary and binary elementwise functions at module import [#1522](https://github.com/IntelPython/dpnp/pull/1522) +* Redesigned trigonometric and hyperbolic functions through pybind11 extension of OneMKL calls where possible or leveraging on `dpctl.tensor` implementation [#1545](https://github.com/IntelPython/dpnp/pull/1545) +* Added `dpnp.signbit` and `dpnp.proj` functions implemented through pybind11 extension of OneMKL calls where possible or leveraging on `dpctl.tensor` implementation [#1535](https://github.com/IntelPython/dpnp/pull/1535) +* Redesigned `dpnp.round` and `dpnp.around` functions through pybind11 extension of OneMKL calls where possible or leveraging on `dpctl.tensor` implementation [#1520](https://github.com/IntelPython/dpnp/pull/1520) +* Redesigned `dpnp.sign` and `dpnp.negative` functions through pybind11 extension of OneMKL calls where possible or leveraging on `dpctl.tensor` implementation [#1523](https://github.com/IntelPython/dpnp/pull/1523) +* Redesigned `dpnp.conjugate` and `dpnp.conj` functions through pybind11 extension of OneMKL calls where possible or leveraging on `dpctl.tensor` implementation [#1519](https://github.com/IntelPython/dpnp/pull/1519) +* Redesigned `dpnp.ceil`, `dpnp.floor` and `dpnp.trunc` functions through pybind11 extension of OneMKL calls where possible or leveraging on `dpctl.tensor` implementation [#1518](https://github.com/IntelPython/dpnp/pull/1518) +* Redesigned `dpnp.remainder` and `dpnp.mod` functions through pybind11 extension of OneMKL calls where possible or leveraging on `dpctl.tensor` implementation [#1515](https://github.com/IntelPython/dpnp/pull/1515) +* Redesigned `dpnp.power` function through pybind11 extension of OneMKL calls where possible or leveraging on `dpctl.tensor` implementation [#1476](https://github.com/IntelPython/dpnp/pull/1476) +* Leveraged `dpctl.tensor` implementation for `dpnp.put` function [#1529](https://github.com/IntelPython/dpnp/pull/1529) +* Leveraged `dpctl.tensor` implementation for `dpnp.roll` and `dpnp.rollaxis` functions [#1517](https://github.com/IntelPython/dpnp/pull/1517) +* Leveraged `dpctl.tensor` implementation for `dpnp.copy` function [#1540](https://github.com/IntelPython/dpnp/pull/1540) +* Leveraged `dpctl.tensor` implementation for `dpnp.expand_dims` and `dpnp.swapaxes` functions [#1532](https://github.com/IntelPython/dpnp/pull/1532) +* Leveraged `dpctl.tensor` implementation for bitwise operations [#1508](https://github.com/IntelPython/dpnp/pull/1508) +* Leveraged `dpctl.tensor` implementation for `dpnp.all` and `dpnp.any` functions [#1512](https://github.com/IntelPython/dpnp/pull/1512) +* Leveraged `dpctl.tensor` implementation for `dpnp.stack` function [#1509](https://github.com/IntelPython/dpnp/pull/1509) +* Leveraged `dpctl.tensor` implementation for `dpnp.concatenate` function [#1507](https://github.com/IntelPython/dpnp/pull/1507) +* Leveraged `dpctl.tensor` implementation for `dpnp.isnan`, `dpnp.isinf` and `dpnp.isfinite` functions [#1504](https://github.com/IntelPython/dpnp/pull/1504) +* Leveraged `dpctl.tensor` implementation for `dpnp.take` function [#1492](https://github.com/IntelPython/dpnp/pull/1492) +* Refreshed API References block in the documentation [#1490](https://github.com/IntelPython/dpnp/pull/1490) +* Refreshed documentation to reflect an actual product behavior [#1485](https://github.com/IntelPython/dpnp/pull/1485) +* Upgraded the build flow to use newer `pybind11=2.11.1` version [#1510](https://github.com/IntelPython/dpnp/pull/1510) +* Updated pre-commit hooks to run with `flake8=6.1.0` and `black=23.7.0` [#1505](https://github.com/IntelPython/dpnp/pull/1505) +* Pinned DPC++ and OneMKL versions to `2023.2`` release [#1496](https://github.com/IntelPython/dpnp/pull/1496) +* Added a specialized kernel for F-contiguous arrays to `dpnp.sum` with `axis=1` [#1489](https://github.com/IntelPython/dpnp/pull/1489) +* Removed a workaround to Klockwork since it is not used anymore due to transition to Coverity tool [#1493](https://github.com/IntelPython/dpnp/pull/1493) + ### Fixed +* Resolved `Logically dead code` issue addressed by Coverity scan [#1541](https://github.com/IntelPython/dpnp/pull/1541) +* Resolved `Arguments in wrong order` issue addressed by Coverity scan [#1513](https://github.com/IntelPython/dpnp/pull/1513) +* Resolved `Pointer to local outside scope` issue addressed by Coverity scan [#1514](https://github.com/IntelPython/dpnp/pull/1514) +* Fixed assigning a value to potentially none-valued dictionary coverage generation script [#1511](https://github.com/IntelPython/dpnp/pull/1511) +* Resolved issues with running `dpnp.allclose` function on a device without fp64 support [#1536](https://github.com/IntelPython/dpnp/pull/1536) +* Resolved issues with running FFT functions on a device without fp64 support [#1524](https://github.com/IntelPython/dpnp/pull/1524) +* Resolved issues with running mathematical functions on a device without fp64 support [#1502](https://github.com/IntelPython/dpnp/pull/1502) +* Resolved issues with running random functions on a device without fp64 support [#1498](https://github.com/IntelPython/dpnp/pull/1498) +* Resolved issues with running statistics functions on a device without fp64 support [#1494](https://github.com/IntelPython/dpnp/pull/1494) + ## [0.12.1] - 07/18/2023 ### Added @@ -29,8 +81,8 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Pinned to `dpctl>=0.14.5` as host and run dependencies [#1481](https://github.com/IntelPython/dpnp/pull/1481) * Pinned dependent `cython` package to a version less than `3.0.0` [#1480](https://github.com/IntelPython/dpnp/pull/1480) * Added a specialized kernel for `dpnp.sum` with `axis=0` as a pybind11 extension [#1479](https://github.com/IntelPython/dpnp/pull/1479) -* Redesigned `dpnp.square` function through pybind11 extension of OneMKL call where possible or Leveraging on `dpctl.tensor` implementation [#1473](https://github.com/IntelPython/dpnp/pull/1473) -* Redesigned `dpnp.cos` and `dpnp.sin` functions through pybind11 extension of OneMKL calls where possible or Leveraging on `dpctl.tensor` implementation [#1471](https://github.com/IntelPython/dpnp/pull/1471) +* Redesigned `dpnp.square` function through pybind11 extension of OneMKL call where possible or leveraging on `dpctl.tensor` implementation [#1473](https://github.com/IntelPython/dpnp/pull/1473) +* Redesigned `dpnp.cos` and `dpnp.sin` functions through pybind11 extension of OneMKL calls where possible or leveraging on `dpctl.tensor` implementation [#1471](https://github.com/IntelPython/dpnp/pull/1471) * Redesigned `dpnp.sqrt` function through pybind11 extension of OneMKL call where possible or leveraging on `dpctl.tensor` implementation [#1470](https://github.com/IntelPython/dpnp/pull/1470) * Redesigned `dpnp.log` function through pybind11 extension of OneMKL call where possible or leveraging on `dpctl.tensor` implementation [#1469](https://github.com/IntelPython/dpnp/pull/1469) * Leveraged `dpctl.tensor` implementation for logical operations [#1464](https://github.com/IntelPython/dpnp/pull/1464) diff --git a/CMakeLists.txt b/CMakeLists.txt index 974d19df6335..b40dbf983a6e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -28,16 +28,16 @@ set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${DPCTL_MODULE_PATH}) find_package(IntelDPCPP REQUIRED) -find_package(TBB REQUIRED PATHS ${CMAKE_SOURCE_DIR}/dpnp/backend/cmake/Modules NO_DEFAULT_PATH) +find_package(TBB REQUIRED) set(MKL_ARCH "intel64") set(MKL_LINK "dynamic") -set(MKL_INTERFACE_FULL "intel_ilp64") set(MKL_THREADING "tbb_thread") -find_package(MKL REQUIRED PATHS ${CMAKE_SOURCE_DIR}/dpnp/backend/cmake/Modules NO_DEFAULT_PATH) +set(MKL_INTERFACE "ilp64") +find_package(MKL REQUIRED) set(ONEDPL_PAR_BACKEND tbb) -find_package(oneDPL REQUIRED PATHS ${CMAKE_SOURCE_DIR}/dpnp/backend/cmake/Modules NO_DEFAULT_PATH) +find_package(oneDPL REQUIRED) include(GNUInstallDirs) diff --git a/README.md b/README.md index 04bc0e0bc3be..b19e902ece1f 100644 --- a/README.md +++ b/README.md @@ -15,11 +15,19 @@ ## Build from source: Ensure you have the following prerequisite packages installed: -- `mkl-devel-dpcpp` +- `cython` +- `cmake >=3.21` - `dpcpp_linux-64` or `dpcpp_win-64` (depending on your OS) +- `dpctl` +- `mkl-devel-dpcpp` - `onedpl-devel` +- `ninja` +- `numpy >=1.19,<1.25a0` +- `python` +- `scikit-build` +- `setuptools` +- `sysroot_linux-64 >=2.28` (only on Linux OS) - `tbb-devel` -- `dpctl` After these steps, `dpnp` can be built in debug mode as follows: diff --git a/conda-recipe/meta.yaml b/conda-recipe/meta.yaml index 339926271d4f..2293cfcdc1f1 100644 --- a/conda-recipe/meta.yaml +++ b/conda-recipe/meta.yaml @@ -14,14 +14,14 @@ requirements: - ninja - git - dpctl >=0.14.5 - - mkl-devel-dpcpp {{ environ.get('MKL_VER', '>=2023.2.0') }} - - onedpl-devel + - mkl-devel-dpcpp {{ environ.get('MKL_VER', '>=2024.0*') }} + - onedpl-devel >=2022.3* - tbb-devel - wheel - scikit-build build: - {{ compiler('cxx') }} - - {{ compiler('dpcpp') }} >=2023.2.0 # [not osx] + - {{ compiler('dpcpp') }} >=2024.0* # [not osx] - sysroot_linux-64 >=2.28 # [linux] run: - python diff --git a/doc/reference/manipulation.rst b/doc/reference/manipulation.rst index 9d8dbd12842f..3f6c0aabfc27 100644 --- a/doc/reference/manipulation.rst +++ b/doc/reference/manipulation.rst @@ -35,6 +35,7 @@ Transpose-like operations :nosignatures: dpnp.moveaxis + dpnp.roll dpnp.rollaxis dpnp.swapaxes dpnp.ndarray.T diff --git a/doc/reference/math.rst b/doc/reference/math.rst index 0d75e08b7d8c..7eb6534fa799 100644 --- a/doc/reference/math.rst +++ b/doc/reference/math.rst @@ -169,6 +169,7 @@ Handling complex numbers dpnp.imag dpnp.conj dpnp.conjugate + dpnp.proj Extrema Finding diff --git a/doc/reference/ufunc.rst b/doc/reference/ufunc.rst index 4e74ac78ca3e..4c1e9f4d5d98 100644 --- a/doc/reference/ufunc.rst +++ b/doc/reference/ufunc.rst @@ -42,6 +42,7 @@ Math operations dpnp.log10 dpnp.expm1 dpnp.log1p + dpnp.proj dpnp.sqrt dpnp.square dpnp.reciprocal diff --git a/dpnp/backend/CMakeLists.txt b/dpnp/backend/CMakeLists.txt index 8aeadc38c376..34f74cf6f46e 100644 --- a/dpnp/backend/CMakeLists.txt +++ b/dpnp/backend/CMakeLists.txt @@ -89,7 +89,7 @@ if(DPNP_GENERATE_COVERAGE) target_link_options(${_trgt} PRIVATE -fprofile-instr-generate -fcoverage-mapping) endif() -target_link_libraries(${_trgt} PUBLIC MKL::MKL_DPCPP) +target_link_libraries(${_trgt} PUBLIC MKL::MKL_SYCL) target_link_libraries(${_trgt} PUBLIC oneDPL) if (UNIX) diff --git a/dpnp/backend/cmake/Modules/MKLConfig.cmake b/dpnp/backend/cmake/Modules/MKLConfig.cmake deleted file mode 100644 index 07290e8ea331..000000000000 --- a/dpnp/backend/cmake/Modules/MKLConfig.cmake +++ /dev/null @@ -1,851 +0,0 @@ -#=============================================================================== -# Copyright 2021-2022 Intel Corporation. -# -# This software and the related documents are Intel copyrighted materials, and -# your use of them is governed by the express license under which they were -# provided to you (License). Unless the License provides otherwise, you may not -# use, modify, copy, publish, distribute, disclose or transmit this software or -# the related documents without Intel's prior written permission. -# -# This software and the related documents are provided as is, with no express -# or implied warranties, other than those that are expressly stated in the -# License. -#=============================================================================== - -#=================================================================== -# CMake Config file for Intel(R) oneAPI Math Kernel Library (oneMKL) -#=================================================================== - -#=============================================================================== -# Input parameters -#================= -#------------- -# Main options -#------------- -# MKL_ROOT: oneMKL root directory (May be required for non-standard install locations. Optional otherwise.) -# Default: use location from MKLROOT environment variable or /../../../ if MKLROOT is not defined -# MKL_ARCH -# Values: ia32 intel64 -# Default: intel64 -# MKL_LINK -# Values: static, dynamic, sdl -# Default: dynamic -# Exceptions:- DPC++ doesn't support sdl -# MKL_THREADING -# Values: sequential, -# intel_thread (Intel OpenMP), -# gnu_thread (GNU OpenMP), -# pgi_thread (PGI OpenMP), -# tbb_thread -# Default: intel_thread -# Exceptions:- DPC++ defaults to tbb, PGI compiler on Windows defaults to pgi_thread -# MKL_INTERFACE (for MKL_ARCH=intel64 only) -# Values: lp64, ilp64 -# GNU or INTEL interface will be selected based on Compiler. -# Default: ilp64 -# MKL_MPI -# Values: intelmpi, mpich, openmpi, msmpi, mshpc -# Default: intelmpi -#----------------------------------- -# Special options (OFF by default) -#----------------------------------- -# ENABLE_BLAS95: Enables BLAS Fortran95 API -# ENABLE_LAPACK95: Enables LAPACK Fortran95 API -# ENABLE_BLACS: Enables cluster BLAS library -# ENABLE_CDFT: Enables cluster DFT library -# ENABLE_CPARDISO: Enables cluster PARDISO functionality -# ENABLE_SCALAPACK: Enables cluster LAPACK library -# ENABLE_OMP_OFFLOAD: Enables OpenMP Offload functionality -# -#================== -# Output parameters -#================== -# MKL_ROOT -# oneMKL root directory. -# MKL_INCLUDE -# Use of target_include_directories() is recommended. -# INTERFACE_INCLUDE_DIRECTORIES property is set on mkl_core and mkl_rt libraries. -# Alternatively, this variable can be used directly (not recommended as per Modern CMake) -# MKL_ENV -# Provides all environment variables based on input parameters. -# Currently useful for mkl_rt linking and BLACS on Windows. -# Must be set as an ENVIRONMENT property. -# Example: -# add_test(NAME mytest COMMAND myexe) -# if(MKL_ENV) -# set_tests_properties(mytest PROPERTIES ENVIRONMENT "${MKL_ENV}") -# endif() -# -# MKL:: -# IMPORTED targets to link MKL libraries individually or when using a custom link-line. -# mkl_core and mkl_rt have INTERFACE_* properties set to them. -# Please refer to Intel(R) oneMKL Link Line Advisor for help with linking. -# -# Below INTERFACE targets provide full link-lines for direct use. -# Example: -# target_link_options( PUBLIC $) -# -# MKL::MKL -# Link line for C and Fortran API -# MKL::MKL_DPCPP -# Link line for DPC++ API -# -# Note: For Device API, library linking is not required. -# Compile options can be added from the INTERFACE_COMPILE_OPTIONS property on MKL::MKL_DPCPP -# Include directories can be added from the INTERFACE_INCLUDE_DIRECTORIES property on MKL::MKL_DPCPP -# -# Note: Output parameters' and targets' availability can change -# based on Input parameters and application project languages. -#=============================================================================== - -function(mkl_message MSG_MODE MSG_TEXT) - if(MSG_MODE STREQUAL "FATAL_ERROR") - message(${MSG_MODE} ${MSG_TEXT}) - else() - if(NOT MKL_FIND_QUIETLY) - message(${MSG_MODE} ${MSG_TEXT}) - endif() - endif() -endfunction() - -if(${CMAKE_VERSION} VERSION_LESS "3.13") - mkl_message(FATAL_ERROR "The minimum supported CMake version is 3.13. You are running version ${CMAKE_VERSION}") -endif() - -include_guard() -include(FindPackageHandleStandardArgs) - -if(NOT MKL_LIBRARIES) - -# Set CMake policies for well-defined behavior across CMake versions -cmake_policy(SET CMP0011 NEW) -cmake_policy(SET CMP0057 NEW) - -# Project Languages -get_property(languages GLOBAL PROPERTY ENABLED_LANGUAGES) -list(APPEND MKL_LANGS C CXX Fortran) -foreach(lang ${languages}) - if(${lang} IN_LIST MKL_LANGS) - list(APPEND CURR_LANGS ${lang}) - endif() -endforeach() -list(REMOVE_DUPLICATES CURR_LANGS) - -option(ENABLE_BLAS95 "Enables BLAS Fortran95 API" OFF) -option(ENABLE_LAPACK95 "Enables LAPACK Fortran95 API" OFF) -option(ENABLE_BLACS "Enables cluster BLAS library" OFF) -option(ENABLE_CDFT "Enables cluster DFT library" OFF) -option(ENABLE_CPARDISO "Enables cluster PARDISO functionality" OFF) -option(ENABLE_SCALAPACK "Enables cluster LAPACK library" OFF) -option(ENABLE_OMP_OFFLOAD "Enables OpenMP Offload functionality" OFF) - -# Use MPI if any of these are enabled -if(ENABLE_BLACS OR ENABLE_CDFT OR ENABLE_SCALAPACK OR ENABLE_CPARDISO) - set(USE_MPI ON) -endif() - -# Check Parameters -function(define_param TARGET_PARAM DEFAULT_PARAM SUPPORTED_LIST) - if(NOT DEFINED ${TARGET_PARAM} AND NOT DEFINED ${DEFAULT_PARAM}) - mkl_message(STATUS "${TARGET_PARAM}: Undefined") - elseif(NOT DEFINED ${TARGET_PARAM} AND DEFINED ${DEFAULT_PARAM}) - set(${TARGET_PARAM} "${${DEFAULT_PARAM}}" CACHE STRING "Choose ${TARGET_PARAM} options are: ${${SUPPORTED_LIST}}") - foreach(opt ${${DEFAULT_PARAM}}) - set(STR_LIST "${STR_LIST} ${opt}") - endforeach() - mkl_message(STATUS "${TARGET_PARAM}: None, set to `${STR_LIST}` by default") - elseif(${SUPPORTED_LIST}) - set(ITEM_FOUND 1) - foreach(opt ${${TARGET_PARAM}}) - if(NOT ${opt} IN_LIST ${SUPPORTED_LIST}) - set(ITEM_FOUND 0) - endif() - endforeach() - if(ITEM_FOUND EQUAL 0) - foreach(opt ${${SUPPORTED_LIST}}) - set(STR_LIST "${STR_LIST} ${opt}") - endforeach() - mkl_message(FATAL_ERROR "Invalid ${TARGET_PARAM} `${${TARGET_PARAM}}`, options are: ${STR_LIST}") - else() - mkl_message(STATUS "${TARGET_PARAM}: ${${TARGET_PARAM}}") - endif() - else() - mkl_message(STATUS "${TARGET_PARAM}: ${${TARGET_PARAM}}") - endif() -endfunction() - -#================ -# Compiler checks -#================ - -if(CMAKE_C_COMPILER) - get_filename_component(C_COMPILER_NAME ${CMAKE_C_COMPILER} NAME) -endif() -if(CMAKE_CXX_COMPILER) - get_filename_component(CXX_COMPILER_NAME ${CMAKE_CXX_COMPILER} NAME) -endif() -if(CMAKE_Fortran_COMPILER) - get_filename_component(Fortran_COMPILER_NAME ${CMAKE_Fortran_COMPILER} NAME) -endif() - -# Determine Compiler Family -if(CXX_COMPILER_NAME STREQUAL "dpcpp" OR CXX_COMPILER_NAME STREQUAL "dpcpp.exe" - OR CXX_COMPILER_NAME STREQUAL "icpx" OR CXX_COMPILER_NAME STREQUAL "icx.exe") - set(DPCPP_COMPILER ON) -endif() -if(C_COMPILER_NAME MATCHES "^clang") - set(CLANG_COMPILER ON) -endif() -if(CMAKE_C_COMPILER_ID STREQUAL "PGI" OR CMAKE_Fortran_COMPILER_ID STREQUAL "PGI") - set(PGI_COMPILER ON) -elseif(CMAKE_C_COMPILER_ID STREQUAL "Intel" OR CMAKE_Fortran_COMPILER_ID STREQUAL "Intel" - OR CMAKE_C_COMPILER_ID STREQUAL "IntelLLVM" OR CMAKE_Fortran_COMPILER_ID STREQUAL "IntelLLVM") - set(INTEL_COMPILER ON) -else() - if(CMAKE_C_COMPILER_ID STREQUAL "GNU") - set(GNU_C_COMPILER ON) - endif() - if(CMAKE_Fortran_COMPILER_ID STREQUAL "GNU") - set(GNU_Fortran_COMPILER ON) - endif() -endif() - -if(USE_MPI AND (C_COMPILER_NAME MATCHES "^mpi" OR Fortran_COMPILER_NAME MATCHES "^mpi")) - set(USE_MPI_SCRIPT ON) -endif() - -#================ - -#================ -# System-specific -#================ - -# Extensions -if(UNIX) - set(LIB_PREFIX "lib") - set(LIB_EXT ".a") - set(DLL_EXT ".so") - if(APPLE) - set(DLL_EXT ".dylib") - endif() - set(LINK_PREFIX "-l") - set(LINK_SUFFIX "") -else() - set(LIB_PREFIX "") - set(LIB_EXT ".lib") - set(DLL_EXT "_dll.lib") - set(LINK_PREFIX "") - set(LINK_SUFFIX ".lib") -endif() - -# Set target system architecture -set(DEFAULT_MKL_ARCH intel64) -if(DPCPP_COMPILER OR PGI_COMPILER OR ENABLE_OMP_OFFLOAD OR USE_MPI) - set(MKL_ARCH_LIST intel64) -else() - set(MKL_ARCH_LIST ia32 intel64) -endif() -define_param(MKL_ARCH DEFAULT_MKL_ARCH MKL_ARCH_LIST) - -#================ - -#========== -# Setup MKL -#========== - -# Set MKL_ROOT directory -if(NOT DEFINED MKL_ROOT) - if(DEFINED ENV{MKLROOT}) - set(MKL_ROOT $ENV{MKLROOT}) - else() - get_filename_component(MKL_CMAKE_PATH "${CMAKE_CURRENT_LIST_DIR}" REALPATH) - get_filename_component(MKL_ROOT "${MKL_CMAKE_PATH}/../../../" ABSOLUTE) - mkl_message(STATUS "MKL_ROOT ${MKL_ROOT}") - endif() -endif() -string(REPLACE "\\" "/" MKL_ROOT ${MKL_ROOT}) - -# Define MKL_LINK -set(DEFAULT_MKL_LINK dynamic) -if(DPCPP_COMPILER OR USE_MPI) - set(MKL_LINK_LIST static dynamic) -else() - set(MKL_LINK_LIST static dynamic sdl) -endif() -define_param(MKL_LINK DEFAULT_MKL_LINK MKL_LINK_LIST) - -# Define MKL_INTERFACE -if(MKL_ARCH STREQUAL "intel64") - set(IFACE_TYPE intel) - if(GNU_Fortran_COMPILER) - set(IFACE_TYPE gf) - endif() - if(DPCPP_COMPILER) - if(MKL_INTERFACE) - set(MKL_INTERFACE_FULL intel_${MKL_INTERFACE}) - endif() - set(DEFAULT_MKL_INTERFACE intel_ilp64) - set(MKL_INTERFACE_LIST intel_ilp64) - else() - if(MKL_INTERFACE) - set(MKL_INTERFACE_FULL ${IFACE_TYPE}_${MKL_INTERFACE}) - endif() - set(DEFAULT_MKL_INTERFACE ${IFACE_TYPE}_ilp64) - set(MKL_INTERFACE_LIST ${IFACE_TYPE}_ilp64 ${IFACE_TYPE}_lp64) - endif() - define_param(MKL_INTERFACE_FULL DEFAULT_MKL_INTERFACE MKL_INTERFACE_LIST) -else() - if(WIN32) - set(MKL_INTERFACE_FULL intel_c) - elseif(NOT APPLE) - if(GNU_Fortran_COMPILER) - set(MKL_INTERFACE_FULL gf) - else() - set(MKL_INTERFACE_FULL intel) - endif() - else() - mkl_message(FATAL_ERROR "OSX does not support MKL_ARCH ia32.") - endif() -endif() -if(MKL_INTERFACE_FULL MATCHES "ilp64") - set(MKL_INTERFACE "ilp64") -else() - set(MKL_INTERFACE "lp64") -endif() - -# Define MKL headers -find_path(MKL_H mkl.h - HINTS ${MKL_ROOT} - PATH_SUFFIXES include) -list(APPEND MKL_INCLUDE ${MKL_H}) - -# Add pre-built F95 Interface Modules -if(INTEL_COMPILER AND (ENABLE_BLAS95 OR ENABLE_LAPACK95)) - if(MKL_ARCH STREQUAL "intel64") - list(APPEND MKL_INCLUDE "${MKL_ROOT}/include/${MKL_ARCH}/${MKL_INTERFACE}") - else() - list(APPEND MKL_INCLUDE "${MKL_ROOT}/include/${MKL_ARCH}") - endif() -endif() - -# Define MKL_THREADING -# All APIs support sequential threading -set(MKL_THREADING_LIST "sequential" "intel_thread" "tbb_thread") -set(DEFAULT_MKL_THREADING intel_thread) -# DPC++ API supports TBB threading, but not OpenMP threading -if(DPCPP_COMPILER) - set(DEFAULT_MKL_THREADING tbb_thread) - list(REMOVE_ITEM MKL_THREADING_LIST intel_thread) -# C, Fortran API -elseif(PGI_COMPILER) - # PGI compiler supports PGI OpenMP threading, additionally - list(APPEND MKL_THREADING_LIST pgi_thread) - # PGI compiler does not support TBB threading - list(REMOVE_ITEM MKL_THREADING_LIST tbb_thread) - if(WIN32) - # PGI 19.10 and 20.1 on Windows, do not support Intel OpenMP threading - list(REMOVE_ITEM MKL_THREADING_LIST intel_thread) - set(DEFAULT_MKL_THREADING pgi_thread) - endif() -elseif(GNU_C_COMPILER OR GNU_Fortran_COMPILER OR CLANG_COMPILER) - list(APPEND MKL_THREADING_LIST gnu_thread) -else() - # Intel and Microsoft compilers - # Nothing to do, only for completeness -endif() -define_param(MKL_THREADING DEFAULT_MKL_THREADING MKL_THREADING_LIST) - -# Define MKL_MPI -set(DEFAULT_MKL_MPI intelmpi) -if(UNIX) - if(APPLE) - # Override defaults for OSX - set(DEFAULT_MKL_MPI mpich) - set(MKL_MPI_LIST mpich) - else() - set(MKL_MPI_LIST intelmpi openmpi mpich mpich2) - endif() -else() - # Windows - set(MKL_MPI_LIST intelmpi mshpc msmpi) -endif() -define_param(MKL_MPI DEFAULT_MKL_MPI MKL_MPI_LIST) -# MSMPI is now called MSHPC. MSMPI option exists for backward compatibility. -if(MKL_MPI STREQUAL "mshpc") - set(MKL_MPI msmpi) -endif() -find_package_handle_standard_args(MKL REQUIRED_VARS MKL_MPI) - -# Checkpoint - Verify if required options are defined -find_package_handle_standard_args(MKL REQUIRED_VARS MKL_ROOT MKL_ARCH MKL_INCLUDE MKL_LINK MKL_THREADING MKL_INTERFACE_FULL) - -# Provides a list of IMPORTED targets for the project -if(NOT DEFINED MKL_IMPORTED_TARGETS) - set(MKL_IMPORTED_TARGETS "") -endif() - -# Clear temporary variables -set(MKL_C_COPT "") -set(MKL_F_COPT "") -set(MKL_SDL_COPT "") -set(MKL_CXX_COPT "") -set(MKL_DPCPP_COPT "") -set(MKL_DPCPP_LOPT "") -set(MKL_OFFLOAD_COPT "") -set(MKL_OFFLOAD_LOPT "") - -set(MKL_SUPP_LINK "") # Other link options. Usually at the end of the link-line. -set(MKL_LINK_LINE) # For MPI only -set(MKL_ENV_PATH "") # Temporary variable to work with PATH -set(MKL_ENV "") # Exported environment variables - -# Modify PATH variable to make it CMake-friendly -set(OLD_PATH $ENV{PATH}) -string(REPLACE ";" "\;" OLD_PATH "${OLD_PATH}") - -# Compiler options -if(GNU_C_COMPILER OR GNU_Fortran_COMPILER) - if(MKL_ARCH STREQUAL "ia32") - list(APPEND MKL_C_COPT -m32) - list(APPEND MKL_F_COPT -m32) - else() - list(APPEND MKL_C_COPT -m64) - list(APPEND MKL_F_COPT -m64) - endif() -endif() - -# Additonal compiler & linker options -if(CXX_COMPILER_NAME STREQUAL "icpx" OR CXX_COMPILER_NAME STREQUAL "icx.exe") - list(APPEND MKL_DPCPP_COPT "-fsycl") - list(APPEND MKL_DPCPP_LOPT "-fsycl") -endif() -if(DPCPP_COMPILER OR ENABLE_OMP_OFFLOAD) - if(MKL_LINK STREQUAL "static") - list(APPEND MKL_DPCPP_LOPT "-fsycl-device-code-split=per_kernel") - list(APPEND MKL_OFFLOAD_LOPT "-fsycl-device-code-split=per_kernel") - endif() -endif() - -# For OpenMP Offload -if(ENABLE_OMP_OFFLOAD) - if(WIN32) - if(OPENMP_VERSION VERSION_GREATER_EQUAL "5.1") - if("Fortran" IN_LIST CURR_LANGS) - list(APPEND MKL_OFFLOAD_COPT -Qiopenmp -Qopenmp-targets:spir64 -DONEMKL_USE_OPENMP_VERSION=202011) - else() - list(APPEND MKL_OFFLOAD_COPT -Qiopenmp -Qopenmp-targets:spir64 -Qopenmp-version:51 -DONEMKL_USE_OPENMP_VERSION=202011) - endif() - else() - list(APPEND MKL_OFFLOAD_COPT -Qiopenmp -Qopenmp-targets:spir64) - endif() - # -MD and -MDd are manually added here because offload functionality uses DPC++ runtime. - if(CMAKE_BUILD_TYPE MATCHES "Debug|DebInfo") - list(APPEND MKL_OFFLOAD_COPT -MDd) - else() - list(APPEND MKL_OFFLOAD_COPT -MD) - endif() - list(APPEND MKL_OFFLOAD_LOPT -Qiopenmp -Qopenmp-targets:spir64 -fsycl) - set(SKIP_LIBPATH ON) - else() - if(OPENMP_VERSION VERSION_GREATER_EQUAL "5.1") - if("Fortran" IN_LIST CURR_LANGS) - list(APPEND MKL_OFFLOAD_COPT -fiopenmp -fopenmp-targets=spir64 -DONEMKL_USE_OPENMP_VERSION=202011) - else() - list(APPEND MKL_OFFLOAD_COPT -fiopenmp -fopenmp-targets=spir64 -fopenmp-version=51 -DONEMKL_USE_OPENMP_VERSION=202011) - endif() - else () - list(APPEND MKL_OFFLOAD_COPT -fiopenmp -fopenmp-targets=spir64) - endif() - list(APPEND MKL_OFFLOAD_LOPT -fiopenmp -fopenmp-targets=spir64 -fsycl) - if(APPLE) - list(APPEND MKL_SUPP_LINK -lc++) - else() - list(APPEND MKL_SUPP_LINK -lstdc++) - endif() - endif() -endif() - -# For selected Interface -if(MKL_INTERFACE_FULL) - if(MKL_ARCH STREQUAL "ia32") - if(GNU_Fortran_COMPILER) - set(MKL_SDL_IFACE_ENV "GNU") - endif() - else() - if(GNU_Fortran_COMPILER) - set(MKL_SDL_IFACE_ENV "GNU,${MKL_INTERFACE}") - else() - set(MKL_SDL_IFACE_ENV "${MKL_INTERFACE}") - endif() - if(MKL_INTERFACE STREQUAL "ilp64") - if("Fortran" IN_LIST CURR_LANGS) - if(INTEL_COMPILER) - if(WIN32) - list(APPEND MKL_F_COPT "-4I8") - else() - list(APPEND MKL_F_COPT "-i8") - endif() - elseif(GNU_Fortran_COMPILER) - list(APPEND MKL_F_COPT "-fdefault-integer-8") - elseif(PGI_COMPILER) - list(APPEND MKL_F_COPT "-i8") - endif() - endif() - list(INSERT MKL_C_COPT 0 "-DMKL_ILP64") - list(INSERT MKL_SDL_COPT 0 "-DMKL_ILP64") - list(INSERT MKL_CXX_COPT 0 "-DMKL_ILP64") - list(INSERT MKL_OFFLOAD_COPT 0 "-DMKL_ILP64") - else() - # lp64 - endif() - endif() - if(MKL_SDL_IFACE_ENV) - string(TOUPPER ${MKL_SDL_IFACE_ENV} MKL_SDL_IFACE_ENV) - endif() -endif() # MKL_INTERFACE_FULL - -# All MKL Libraries -if(WIN32 AND CMAKE_BUILD_TYPE MATCHES "Debug|DebInfo") - set(MKL_SYCL mkl_sycld) -else() - set(MKL_SYCL mkl_sycl) -endif() -set(MKL_IFACE_LIB mkl_${MKL_INTERFACE_FULL}) -set(MKL_CORE mkl_core) -if(WIN32 AND CMAKE_BUILD_TYPE MATCHES "Debug|DebInfo" AND MKL_THREADING STREQUAL "tbb_thread") - set(MKL_THREAD mkl_tbb_threadd) -else() - set(MKL_THREAD mkl_${MKL_THREADING}) -endif() -set(MKL_SDL mkl_rt) -if(MKL_ARCH STREQUAL "ia32") - set(MKL_BLAS95 mkl_blas95) - set(MKL_LAPACK95 mkl_lapack95) -else() - set(MKL_BLAS95 mkl_blas95_${MKL_INTERFACE}) - set(MKL_LAPACK95 mkl_lapack95_${MKL_INTERFACE}) -endif() -# BLACS -set(MKL_BLACS mkl_blacs_${MKL_MPI}_${MKL_INTERFACE}) -if(UNIX AND NOT APPLE AND MKL_MPI MATCHES "mpich") - # MPICH is compatible with INTELMPI Wrappers on Linux - set(MKL_BLACS mkl_blacs_intelmpi_${MKL_INTERFACE}) -endif() -if(WIN32) - if(MKL_MPI STREQUAL "msmpi") - if("Fortran" IN_LIST CURR_LANGS) - list(APPEND MKL_SUPP_LINK "msmpifec.lib") - endif() - # MSMPI and MSHPC are supported with the same BLACS library - set(MKL_BLACS mkl_blacs_msmpi_${MKL_INTERFACE}) - if(NOT MKL_LINK STREQUAL "static") - set(MKL_BLACS mkl_blacs_${MKL_INTERFACE}) - set(MKL_BLACS_ENV MSMPI) - endif() - elseif(MKL_MPI STREQUAL "intelmpi" AND NOT MKL_LINK STREQUAL "static") - set(MKL_BLACS mkl_blacs_${MKL_INTERFACE}) - set(MKL_BLACS_ENV INTELMPI) - endif() -endif() -# CDFT & SCALAPACK -set(MKL_CDFT mkl_cdft_core) -set(MKL_SCALAPACK mkl_scalapack_${MKL_INTERFACE}) - - -if (UNIX) - if(NOT APPLE) - if(MKL_LINK STREQUAL "static") - set(START_GROUP "-Wl,--start-group") - set(END_GROUP "-Wl,--end-group") - if(DPCPP_COMPILER OR ENABLE_OMP_OFFLOAD) - set(EXPORT_DYNAMIC "-Wl,-export-dynamic") - endif() - elseif(MKL_LINK STREQUAL "dynamic") - set(MKL_RPATH "-Wl,-rpath=$") - if((GNU_Fortran_COMPILER OR PGI_COMPILER) AND "Fortran" IN_LIST CURR_LANGS) - set(NO_AS_NEEDED -Wl,--no-as-needed) - endif() - else() - set(MKL_RPATH "-Wl,-rpath=$") - endif() - endif() -endif() - -# Create a list of requested libraries, based on input options (MKL_LIBRARIES) -# Create full link-line in MKL_LINK_LINE -list(APPEND MKL_LINK_LINE $,${MKL_OFFLOAD_LOPT},> - $,${MKL_DPCPP_LOPT},> ${EXPORT_DYNAMIC} ${NO_AS_NEEDED} ${MKL_RPATH}) -if(ENABLE_BLAS95) - list(APPEND MKL_LIBRARIES ${MKL_BLAS95}) - list(APPEND MKL_LINK_LINE MKL::${MKL_BLAS95}) -endif() -if(ENABLE_LAPACK95) - list(APPEND MKL_LIBRARIES ${MKL_LAPACK95}) - list(APPEND MKL_LINK_LINE MKL::${MKL_LAPACK95}) -endif() -if(ENABLE_SCALAPACK) - list(APPEND MKL_LIBRARIES ${MKL_SCALAPACK}) - list(APPEND MKL_LINK_LINE MKL::${MKL_SCALAPACK}) -endif() -if(DPCPP_COMPILER OR (ENABLE_OMP_OFFLOAD AND NOT MKL_LINK STREQUAL "sdl")) - list(APPEND MKL_LIBRARIES ${MKL_SYCL}) - list(APPEND MKL_LINK_LINE MKL::${MKL_SYCL}) -endif() -list(APPEND MKL_LINK_LINE ${START_GROUP}) -if(ENABLE_CDFT) - list(APPEND MKL_LIBRARIES ${MKL_CDFT}) - list(APPEND MKL_LINK_LINE MKL::${MKL_CDFT}) -endif() -if(MKL_LINK STREQUAL "sdl") - list(APPEND MKL_LIBRARIES ${MKL_SDL}) - list(APPEND MKL_LINK_LINE MKL::${MKL_SDL}) -else() - list(APPEND MKL_LIBRARIES ${MKL_IFACE_LIB} ${MKL_THREAD} ${MKL_CORE}) - list(APPEND MKL_LINK_LINE MKL::${MKL_IFACE_LIB} MKL::${MKL_THREAD} MKL::${MKL_CORE}) -endif() -if(USE_MPI) - list(APPEND MKL_LIBRARIES ${MKL_BLACS}) - list(APPEND MKL_LINK_LINE MKL::${MKL_BLACS}) -endif() -list(APPEND MKL_LINK_LINE ${END_GROUP}) - -# Find all requested libraries -foreach(lib ${MKL_LIBRARIES}) - unset(${lib}_file CACHE) - if(MKL_LINK STREQUAL "static" AND NOT ${lib} STREQUAL ${MKL_SDL}) - find_library(${lib}_file ${LIB_PREFIX}${lib}${LIB_EXT} - PATHS ${MKL_ROOT} - PATH_SUFFIXES "lib" "lib/${MKL_ARCH}") - add_library(MKL::${lib} STATIC IMPORTED) - else() - find_library(${lib}_file NAMES ${LIB_PREFIX}${lib}${DLL_EXT} ${lib} - PATHS ${MKL_ROOT} - PATH_SUFFIXES "lib" "lib/${MKL_ARCH}") - add_library(MKL::${lib} SHARED IMPORTED) - endif() - find_package_handle_standard_args(MKL REQUIRED_VARS ${lib}_file) - # CMP0111, implemented in CMake 3.20+ requires a shared library target on Windows - # to be defined with IMPLIB and LOCATION property. - # It also requires a static library target to be defined with LOCATION property. - # Setting the policy to OLD usage, using cmake_policy() does not work as of 3.20.0, hence the if-else below. - if(WIN32 AND NOT MKL_LINK STREQUAL "static") - set_target_properties(MKL::${lib} PROPERTIES IMPORTED_IMPLIB "${${lib}_file}") - # Find corresponding DLL - set(MKL_DLL_GLOB ${lib}.*.dll) - file(GLOB MKL_DLL_FILE "${MKL_ROOT}/redist/${MKL_ARCH}/${MKL_DLL_GLOB}" - "${MKL_ROOT}/../redist/${MKL_ARCH}/${MKL_DLL_GLOB}" - "${MKL_ROOT}/../redist/${MKL_ARCH}/mkl/${MKL_DLL_GLOB}" - "${MKL_ROOT}/bin/${MKL_DLL_GLOB}") - if(NOT ${lib} STREQUAL ${MKL_IFACE_LIB} AND NOT ${lib} STREQUAL ${MKL_BLAS95} AND NOT ${lib} STREQUAL ${MKL_LAPACK95}) # Windows IFACE libs are static only - list(LENGTH MKL_DLL_FILE MKL_DLL_FILE_LEN) - if(MKL_DLL_FILE_LEN) - # in case multiple versions of the same dll are found, select the highest version - list(SORT MKL_DLL_FILE) - list(REVERSE MKL_DLL_FILE) - list(GET MKL_DLL_FILE 0 MKL_DLL_FILE) - - mkl_message(STATUS "Found DLL: ${MKL_DLL_FILE}") - set_target_properties(MKL::${lib} PROPERTIES IMPORTED_LOCATION "${MKL_DLL_FILE}") - else() - mkl_message(FATAL_ERROR "${MKL_DLL_GLOB} not found. MKL_ROOT was '${MKL_ROOT}'. MKL_DLL_FILE is '${MKL_DLL_FILE}'") - endif() - endif() - else() - set_target_properties(MKL::${lib} PROPERTIES IMPORTED_LOCATION "${${lib}_file}") - endif() - list(APPEND MKL_IMPORTED_TARGETS MKL::${lib}) -endforeach() - -# Threading selection -if(MKL_THREADING) - if(MKL_THREADING STREQUAL "tbb_thread") - find_package(TBB REQUIRED CONFIG COMPONENTS tbb) - set(MKL_THREAD_LIB $) - set(MKL_SDL_THREAD_ENV "TBB") - get_property(TBB_LIB TARGET TBB::tbb PROPERTY IMPORTED_LOCATION_RELEASE) - get_filename_component(TBB_LIB_DIR ${TBB_LIB} DIRECTORY) - if(UNIX) - if(CMAKE_SKIP_BUILD_RPATH) - set(TBB_LINK "-L${TBB_LIB_DIR} -ltbb") - else() - set(TBB_LINK "-Wl,-rpath,${TBB_LIB_DIR} -L${TBB_LIB_DIR} -ltbb") - endif() - list(APPEND MKL_SUPP_LINK ${TBB_LINK}) - if(APPLE) - list(APPEND MKL_SUPP_LINK -lc++) - else() - list(APPEND MKL_SUPP_LINK -lstdc++) - endif() - endif() - if(WIN32 OR APPLE) - set(MKL_ENV_PATH ${TBB_LIB_DIR}) - endif() - elseif(MKL_THREADING MATCHES "_thread") - if(MKL_THREADING STREQUAL "pgi_thread") - list(APPEND MKL_SUPP_LINK -mp -pgf90libs) - set(MKL_SDL_THREAD_ENV "PGI") - elseif(MKL_THREADING STREQUAL "gnu_thread") - list(APPEND MKL_SUPP_LINK -lgomp) - set(MKL_SDL_THREAD_ENV "GNU") - else() - # intel_thread - if(UNIX) - set(MKL_OMP_LIB iomp5) - set(LIB_EXT ".so") - if(APPLE) - set(LIB_EXT ".dylib") - endif() - else() - set(MKL_OMP_LIB libiomp5md) - endif() - set(MKL_SDL_THREAD_ENV "INTEL") - set(OMP_LIBNAME ${LIB_PREFIX}${MKL_OMP_LIB}${LIB_EXT}) - - find_library(OMP_LIBRARY ${OMP_LIBNAME} - HINTS $ENV{LIB} $ENV{LIBRARY_PATH} $ENV{MKLROOT} ${MKL_ROOT} ${CMPLR_ROOT} - PATH_SUFFIXES "lib" "lib/${MKL_ARCH}" - "lib/${MKL_ARCH}_lin" "lib/${MKL_ARCH}_win" - "linux/compiler/lib/${MKL_ARCH}" - "linux/compiler/lib/${MKL_ARCH}_lin" - "windows/compiler/lib/${MKL_ARCH}" - "windows/compiler/lib/${MKL_ARCH}_win" - "../compiler/lib/${MKL_ARCH}_lin" "../compiler/lib/${MKL_ARCH}_win" - "../compiler/lib/${MKL_ARCH}" "../compiler/lib" - "../../compiler/latest/linux/compiler/lib/${MKL_ARCH}" - "../../compiler/latest/linux/compiler/lib/${MKL_ARCH}_lin" - "../../compiler/latest/windows/compiler/lib/${MKL_ARCH}" - "../../compiler/latest/windows/compiler/lib/${MKL_ARCH}_win" - "../../compiler/latest/mac/compiler/lib") - if(WIN32) - set(OMP_DLLNAME ${LIB_PREFIX}${MKL_OMP_LIB}.dll) - find_path(OMP_DLL_DIR ${OMP_DLLNAME} - HINTS $ENV{LIB} $ENV{LIBRARY_PATH} $ENV{MKLROOT} ${MKL_ROOT} ${CMPLR_ROOT} - PATH_SUFFIXES "redist/${MKL_ARCH}" - "redist/${MKL_ARCH}_win" "redist/${MKL_ARCH}_win/compiler" - "../redist/${MKL_ARCH}/compiler" "../compiler/lib" - "../../compiler/latest/windows/redist/${MKL_ARCH}_win" - "../../compiler/latest/windows/redist/${MKL_ARCH}_win/compiler" - "../../compiler/latest/windows/compiler/redist/${MKL_ARCH}_win" - "../../compiler/latest/windows/compiler/redist/${MKL_ARCH}_win/compiler") - find_package_handle_standard_args(MKL REQUIRED_VARS OMP_DLL_DIR) - set(MKL_ENV_PATH "${OMP_DLL_DIR}") - endif() - - if(WIN32 AND SKIP_LIBPATH) - # Only for Intel OpenMP Offload - set(OMP_LINK "libiomp5md.lib") - else() - set(OMP_LINK "${OMP_LIBRARY}") - if(CMAKE_C_COMPILER_ID STREQUAL "PGI" OR CMAKE_Fortran_COMPILER_ID STREQUAL "PGI") - # Disable PGI OpenMP runtime for correct work of Intel OpenMP runtime - list(APPEND MKL_SUPP_LINK -nomp) - endif() - endif() - find_package_handle_standard_args(MKL REQUIRED_VARS OMP_LIBRARY OMP_LINK) - set(MKL_THREAD_LIB ${OMP_LINK}) - endif() - else() - # Sequential threading - set(MKL_SDL_THREAD_ENV "SEQUENTIAL") - endif() -endif() # MKL_THREADING - -if (UNIX) - list(APPEND MKL_SUPP_LINK -lm -ldl -lpthread) -endif() - -if(DPCPP_COMPILER OR ENABLE_OMP_OFFLOAD) - if(WIN32) - # Detect sycl library version - if(NOT DEFINED SYCL_LIB_VER_CACHE) - set(SYCL_LIB_VER "") - find_library(SYCL_LIB_DIR ${LIB_PREFIX}sycl${LIB_EXT} - HINTS $ENV{LIB} $ENV{CMPLR_ROOT} - PATH_SUFFIXES "windows/lib") - if(NOT SYCL_LIB_DIR) - foreach(ver RANGE 6 99) - find_library(SYCL_LIB_DIR ${LIB_PREFIX}sycl${ver}${LIB_EXT} - HINTS $ENV{LIB} $ENV{CMPLR_ROOT} - PATH_SUFFIXES "windows/lib") - if(SYCL_LIB_DIR) - set(SYCL_LIB_VER ${ver}) - break() - endif() - endforeach() - endif() - set(SYCL_LIB_VER_CACHE ${SYCL_LIB_VER} CACHE STRING "") - endif() - - if(CMAKE_BUILD_TYPE MATCHES "Debug|DebInfo") - list(APPEND MKL_SUPP_LINK ${LINK_PREFIX}sycl${SYCL_LIB_VER_CACHE}d${LINK_SUFFIX}) - else() - list(APPEND MKL_SUPP_LINK ${LINK_PREFIX}sycl${SYCL_LIB_VER_CACHE}${LINK_SUFFIX}) - endif() - else() - list(APPEND MKL_SUPP_LINK ${LINK_PREFIX}sycl${LINK_SUFFIX}) - endif() - list(APPEND MKL_SUPP_LINK ${LINK_PREFIX}OpenCL${LINK_SUFFIX}) -endif() - -# Setup link types based on input options -set(LINK_TYPES "") - -if(DPCPP_COMPILER) - add_library(MKL::MKL_DPCPP INTERFACE IMPORTED GLOBAL) - target_compile_options(MKL::MKL_DPCPP INTERFACE ${MKL_DPCPP_COPT}) - target_link_libraries(MKL::MKL_DPCPP INTERFACE ${MKL_LINK_LINE} ${MKL_THREAD_LIB} ${MKL_SUPP_LINK}) - list(APPEND LINK_TYPES MKL::MKL_DPCPP) -endif() -# Single target for all C, Fortran link-lines -add_library(MKL::MKL INTERFACE IMPORTED GLOBAL) -target_compile_options(MKL::MKL INTERFACE - $<$,C>:${MKL_C_COPT}> - $<$,Fortran>:${MKL_F_COPT}> - $<$,CXX>:${MKL_CXX_COPT}> - $,${MKL_OFFLOAD_COPT},>) -target_link_libraries(MKL::MKL INTERFACE ${MKL_LINK_LINE} ${MKL_THREAD_LIB} ${MKL_SUPP_LINK}) -list(APPEND LINK_TYPES MKL::MKL) - -foreach(link ${LINK_TYPES}) - # Set properties on all INTERFACE targets - target_include_directories(${link} BEFORE INTERFACE "${MKL_INCLUDE}") - list(APPEND MKL_IMPORTED_TARGETS ${link}) -endforeach(link) # LINK_TYPES - -if(MKL_LINK STREQUAL "sdl") - list(APPEND MKL_ENV "MKL_INTERFACE_LAYER=${MKL_SDL_IFACE_ENV}" "MKL_THREADING_LAYER=${MKL_SDL_THREAD_ENV}") -endif() -if(WIN32 AND NOT MKL_LINK STREQUAL "static") - list(APPEND MKL_ENV "MKL_BLACS_MPI=${MKL_BLACS_ENV}") -endif() - -# Add MKL dynamic libraries if RPATH is not defined on Unix -if(UNIX AND CMAKE_SKIP_BUILD_RPATH) - if(MKL_LINK STREQUAL "sdl") - set(MKL_LIB_DIR $) - else() - set(MKL_LIB_DIR $) - endif() - if(APPLE) - list(APPEND MKL_ENV "DYLD_LIBRARY_PATH=${MKL_LIB_DIR}\;$ENV{DYLD_LIBRARY_PATH}") - else() - list(APPEND MKL_ENV "LD_LIBRARY_PATH=${MKL_LIB_DIR}\;$ENV{LD_LIBRARY_PATH}") - endif() -endif() - -# Add MKL dynamic libraries to PATH on Windows -if(WIN32 AND NOT MKL_LINK STREQUAL "static") - get_filename_component(MKL_DLL_DIR ${MKL_DLL_FILE} DIRECTORY) - set(MKL_ENV_PATH "${MKL_DLL_DIR}\;${MKL_ENV_PATH}") -endif() - -if(MKL_ENV_PATH) - list(APPEND MKL_ENV "PATH=${MKL_ENV_PATH}\;${OLD_PATH}") - if(APPLE) - list(APPEND MKL_ENV "DYLD_LIBRARY_PATH=${MKL_ENV_PATH}\:${OLD_PATH}") - endif() -endif() - -unset(MKL_DLL_FILE) - -endif() # MKL_LIBRARIES diff --git a/dpnp/backend/cmake/Modules/README.md b/dpnp/backend/cmake/Modules/README.md deleted file mode 100644 index 0a6b2df4089c..000000000000 --- a/dpnp/backend/cmake/Modules/README.md +++ /dev/null @@ -1,8 +0,0 @@ -# oneAPI CMake scripts vendored from Intel oneAPI BaseKit 2023.0.0 - -This is done to work around absence of this script in onedpl-devel conda -package. Once it is added, expected 2023.2.0, this vendored package is -to be removed. - -tbb-devel script has been modified to allow it to work correctly in conda -environment. diff --git a/dpnp/backend/cmake/Modules/TBBConfig.cmake b/dpnp/backend/cmake/Modules/TBBConfig.cmake deleted file mode 100644 index 5363c1d3808b..000000000000 --- a/dpnp/backend/cmake/Modules/TBBConfig.cmake +++ /dev/null @@ -1,193 +0,0 @@ -# Copyright (c) 2017-2023 Intel Corporation -# -# Licensed under the Apache License, Version 2.0 (the "License"); -# you may not use this file except in compliance with the License. -# You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. - -# It defines the following variables: -# TBB__FOUND -# TBB_IMPORTED_TARGETS -# -# TBBConfigVersion.cmake defines TBB_VERSION -# -# Initialize to default values -if (NOT TBB_IMPORTED_TARGETS) - set(TBB_IMPORTED_TARGETS "") -endif() - -if (NOT TBB_FIND_COMPONENTS) - set(TBB_FIND_COMPONENTS "tbb;tbbmalloc;tbbmalloc_proxy") - foreach (_tbb_component ${TBB_FIND_COMPONENTS}) - set(TBB_FIND_REQUIRED_${_tbb_component} 1) - endforeach() -endif() - -get_filename_component(_tbb_root "${CMAKE_CURRENT_LIST_DIR}" REALPATH) -get_filename_component(_tbb_root "${_tbb_root}/../../.." ABSOLUTE) - -set(TBB_INTERFACE_VERSION ) - -set(_tbb_bin_version 12) -set(_tbbmalloc_bin_version 2) -set(_tbbmalloc_proxy_bin_version 2) -set(_tbbbind_bin_version 3) - -# Add components with internal dependencies: tbbmalloc_proxy -> tbbmalloc -list(FIND TBB_FIND_COMPONENTS tbbmalloc_proxy _tbbmalloc_proxy_ix) -if (NOT _tbbmalloc_proxy_ix EQUAL -1) - list(APPEND TBB_FIND_COMPONENTS tbbmalloc) - list(REMOVE_DUPLICATES TBB_FIND_COMPONENTS) - set(TBB_FIND_REQUIRED_tbbmalloc ${TBB_FIND_REQUIRED_tbbmalloc_proxy}) -endif() -unset(_tbbmalloc_proxy_ix) - -if (CMAKE_SIZEOF_VOID_P STREQUAL "8") - set(_tbb_subdir intel64/gcc4.8) -else () - set(_tbb_subdir ia32/gcc4.8) -endif() - -if (UNIX) - set(_tbb_lib_ext ".so") - set(_tbb_lib_prefix "lib") - set(_tbb_lib_dir_conda "lib") - set(_bin_version "") -elseif (WIN32) - set(_bin_version "") - set(_tbb_lib_prefix "") - set(_tbb_lib_ext ".dll") - set(_tbb_impllib_ext ".lib") - set(_tbb_lib_dir_conda "bin") - set(_tbb_impllib_dir_conda "lib") -else() - message(FATAL_ERROR "Unsupported platform. Only Unix and Windows are supported.") -endif() - -foreach (_tbb_component ${TBB_FIND_COMPONENTS}) - set(TBB_${_tbb_component}_FOUND 0) - -if(WIN32) - unset(_bin_version) - if (_tbb_component STREQUAL tbb) - set(_bin_version ${_tbb_bin_version}) - endif() -endif() - - if(UNIX) - find_library(_tbb_release_lib - NAMES ${_tbb_lib_prefix}${_tbb_component}${_bin_version}${_tbb_lib_ext} - PATHS ${_tbb_root} - HINTS ENV TBB_ROOT_HINT - PATH_SUFFIXES "${_tbb_lib_dir_conda}" "lib/${_tbb_subdir}") - - else() - find_file(_tbb_release_lib - NAMES ${_tbb_lib_prefix}${_tbb_component}${_bin_version}${_tbb_lib_ext} - PATHS ${_tbb_root} - HINTS ENV TBB_ROOT_HINT - PATH_SUFFIXES "${_tbb_lib_dir_conda}" "lib/${_tbb_subdir}") - - if (EXISTS "${_tbb_release_lib}") - find_library(_tbb_release_impllib - NAMES ${_tbb_lib_prefix}${_tbb_component}${_bin_version}${_tbb_impllib_ext} - PATHS ${_tbb_root} - HINTS ENV TBB_ROOT_HINT - PATH_SUFFIXES "${_tbb_impllib_dir_conda}" "lib/${_tbb_subdir}") - endif() - endif() - - if (NOT TBB_FIND_RELEASE_ONLY) - find_library(_tbb_debug_lib - NAMES ${_tbb_lib_prefix}${_tbb_component}${_bin_version}_debug.${_tbb_lib_ext} - PATHS ${_tbb_root} - HINTS ENV TBB_ROOT_HINT - PATH_SUFFIXES "${_tbb_lib_dir_conda}" "lib/${_tbb_subdir}") - if(WIN32 AND EXISTS "${_tbb_debug_lib}") - find_library(_tbb_debug_impllib - NAMES ${_tbb_lib_prefix}${_tbb_component}${_bin_version}_debug.${_tbb_impllib_ext} - PATHS ${_tbb_root} - HINTS ENV TBB_ROOT_HINT - PATH_SUFFIXES "${_tbb_impllib_dir_conda}" "lib/${_tbb_subdir}") - endif() - endif() - - if (EXISTS "${_tbb_release_lib}" OR EXISTS "${_tbb_debug_lib}") - if (NOT TARGET TBB::${_tbb_component}) - add_library(TBB::${_tbb_component} SHARED IMPORTED) - - find_path(_tbb_include_dir - oneapi/tbb.h - PATHS ${_tbb_root} - PATH_SUFFIXES include - HINTS ENV TBB_ROOT_HINT - ) - -if(WIN32) - set_target_properties( - TBB::${_tbb_component} PROPERTIES - INTERFACE_INCLUDE_DIRECTORIES "${_tbb_include_dir}" - INTERFACE_COMPILE_DEFINITIONS "__TBB_NO_IMPLICIT_LINKAGE=1" - ) -else() - set_target_properties( - TBB::${_tbb_component} PROPERTIES - INTERFACE_INCLUDE_DIRECTORIES "${_tbb_include_dir}" - ) -endif() - unset(_tbb_current_realpath) - unset(_tbb_include_dir) - - if (EXISTS "${_tbb_release_lib}") -if(WIN32) - set_target_properties(TBB::${_tbb_component} PROPERTIES - IMPORTED_LOCATION_RELEASE "${_tbb_release_lib}" - IMPORTED_IMPLIB_RELEASE "${_tbb_release_impllib}") -else() - set_target_properties(TBB::${_tbb_component} PROPERTIES - IMPORTED_LOCATION_RELEASE "${_tbb_release_lib}") -endif() - set_property(TARGET TBB::${_tbb_component} APPEND PROPERTY IMPORTED_CONFIGURATIONS RELEASE) - endif() - - if (EXISTS "${_tbb_debug_lib}") -if(WIN32) - set_target_properties(TBB::${_tbb_component} PROPERTIES - IMPORTED_LOCATION_DEBUG "${_tbb_debug_lib}" - IMPORTED_IMPLIB_DEBUG "${_tbb_debug_impllib}" - ) -else() - set_target_properties(TBB::${_tbb_component} PROPERTIES - IMPORTED_LOCATION_DEBUG "${_tbb_debug_lib}") -endif() - set_property(TARGET TBB::${_tbb_component} APPEND PROPERTY IMPORTED_CONFIGURATIONS DEBUG) - endif() - - # Add internal dependencies for imported targets: TBB::tbbmalloc_proxy -> TBB::tbbmalloc - if (_tbb_component STREQUAL tbbmalloc_proxy) - set_target_properties(TBB::tbbmalloc_proxy PROPERTIES INTERFACE_LINK_LIBRARIES TBB::tbbmalloc) - endif() - endif() - list(APPEND TBB_IMPORTED_TARGETS TBB::${_tbb_component}) - set(TBB_${_tbb_component}_FOUND 1) - elseif (TBB_FIND_REQUIRED AND TBB_FIND_REQUIRED_${_tbb_component}) - message(STATUS "Missed required oneTBB component: ${_tbb_component}") - if (TBB_FIND_RELEASE_ONLY) - message(STATUS " ${_tbb_release_lib} must exist.") - else() - message(STATUS " one or both of:\n ${_tbb_release_lib}\n ${_tbb_debug_lib}\n files must exist.") - endif() - set(TBB_FOUND FALSE) - endif() -endforeach() -list(REMOVE_DUPLICATES TBB_IMPORTED_TARGETS) -unset(_tbb_release_lib) -unset(_tbb_debug_lib) -unset(_tbb_root) diff --git a/dpnp/backend/cmake/Modules/oneDPLConfig.cmake b/dpnp/backend/cmake/Modules/oneDPLConfig.cmake deleted file mode 100644 index 6473d20c69f5..000000000000 --- a/dpnp/backend/cmake/Modules/oneDPLConfig.cmake +++ /dev/null @@ -1,105 +0,0 @@ -##===----------------------------------------------------------------------===## -# -# Copyright (C) Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -# This file incorporates work covered by the following copyright and permission -# notice: -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# -##===----------------------------------------------------------------------===## - -# Installation path: /lib/cmake/oneDPL/ -get_filename_component(_onedpl_root "${CMAKE_CURRENT_LIST_DIR}" REALPATH) -get_filename_component(_onedpl_root "${_onedpl_root}/../../../" ABSOLUTE) - -if (WIN32) - set(_onedpl_headers_subdir windows) -else() - set(_onedpl_headers_subdir linux) -endif() - - -find_path(_onedpl_headers - NAMES oneapi/dpl - PATHS ${_onedpl_root} - HINTS ENV DPL_ROOT_HINT - PATH_SUFFIXES include ${_onedpl_headers_subdir}/include -) - - -if (EXISTS "${_onedpl_headers}") - if (NOT TARGET oneDPL) - include(CheckCXXCompilerFlag) - - add_library(oneDPL INTERFACE IMPORTED) - set_target_properties(oneDPL PROPERTIES INTERFACE_INCLUDE_DIRECTORIES "${_onedpl_headers}") - - if (ONEDPL_PAR_BACKEND AND NOT ONEDPL_PAR_BACKEND MATCHES "^(tbb|openmp|serial)$") - message(STATUS "oneDPL: ONEDPL_PAR_BACKEND=${ONEDPL_PAR_BACKEND} is requested, but not supported, available backends: tbb, openmp, serial") - set(oneDPL_FOUND FALSE) - return() - endif() - - if (NOT ONEDPL_PAR_BACKEND OR ONEDPL_PAR_BACKEND STREQUAL "tbb") # Handle oneTBB backend - if (NOT TBB_FOUND) - find_package(TBB 2021 QUIET COMPONENTS tbb PATHS ${CMAKE_SOURCE_DIR}/dpnp/backend/cmake/Modules NO_DEFAULT_PATH) - endif() - if (NOT TBB_FOUND AND ONEDPL_PAR_BACKEND STREQUAL "tbb") # If oneTBB backend is requested explicitly, but not found. - message(STATUS "oneDPL: ONEDPL_PAR_BACKEND=${ONEDPL_PAR_BACKEND} requested, but not found") - set(oneDPL_FOUND FALSE) - return() - elseif (TBB_FOUND) - set(ONEDPL_PAR_BACKEND tbb) - message(STATUS "oneDPL: ONEDPL_PAR_BACKEND=${ONEDPL_PAR_BACKEND}, disable OpenMP backend") - set_target_properties(oneDPL PROPERTIES INTERFACE_LINK_LIBRARIES TBB::tbb) - set_property(TARGET oneDPL APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS ONEDPL_USE_TBB_BACKEND=1 ONEDPL_USE_OPENMP_BACKEND=0) - endif() - endif() - - if (NOT ONEDPL_PAR_BACKEND OR ONEDPL_PAR_BACKEND STREQUAL "openmp") # Handle OpenMP backend - if (UNIX) - set(_openmp_flag "-fopenmp") - else() - set(_openmp_flag "-Qopenmp") - endif() - - # Some compilers may fail if _openmp_flag is not in CMAKE_REQUIRED_LIBRARIES. - set(_onedpl_saved_required_libs ${CMAKE_REQUIRED_LIBRARIES}) - set(CMAKE_REQUIRED_LIBRARIES ${_openmp_option}) - check_cxx_compiler_flag(${_openmp_flag} _openmp_option) - set(CMAKE_REQUIRED_LIBRARIES ${_onedpl_saved_required_libs}) - unset(_onedpl_saved_required_libs) - - if (NOT _openmp_option AND ONEDPL_PAR_BACKEND STREQUAL "openmp") # If OpenMP backend is requested explicitly, but not supported. - message(STATUS "oneDPL: ONEDPL_PAR_BACKEND=${ONEDPL_PAR_BACKEND} requested, but not supported") - set(oneDPL_FOUND FALSE) - return() - elseif (_openmp_option) - set(ONEDPL_PAR_BACKEND openmp) - message(STATUS "oneDPL: ONEDPL_PAR_BACKEND=${ONEDPL_PAR_BACKEND}, disable oneTBB backend") - set_target_properties(oneDPL PROPERTIES INTERFACE_COMPILE_OPTIONS ${_openmp_flag}) - set_target_properties(oneDPL PROPERTIES INTERFACE_LINK_LIBRARIES ${_openmp_flag}) - set_property(TARGET oneDPL APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS ONEDPL_USE_TBB_BACKEND=0 ONEDPL_USE_OPENMP_BACKEND=1) - endif() - endif() - - if (NOT ONEDPL_PAR_BACKEND OR ONEDPL_PAR_BACKEND STREQUAL "serial") - set(ONEDPL_PAR_BACKEND serial) - message(STATUS "oneDPL: ONEDPL_PAR_BACKEND=${ONEDPL_PAR_BACKEND}, disable oneTBB and OpenMP backends") - set_property(TARGET oneDPL APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS ONEDPL_USE_TBB_BACKEND=0 ONEDPL_USE_OPENMP_BACKEND=0) - endif() - - check_cxx_compiler_flag("-fsycl" _fsycl_option) - if (NOT _fsycl_option) - message(STATUS "oneDPL: -fsycl is not supported by current compiler, set ONEDPL_USE_DPCPP_BACKEND=0") - set_property(TARGET oneDPL APPEND PROPERTY INTERFACE_COMPILE_DEFINITIONS ONEDPL_USE_DPCPP_BACKEND=0) - endif() - endif() -else() - message(STATUS "oneDPL: headers do not exist ${_onedpl_headers}") - set(oneDPL_FOUND FALSE) -endif() diff --git a/dpnp/backend/extensions/lapack/CMakeLists.txt b/dpnp/backend/extensions/lapack/CMakeLists.txt index e54de4068c01..9ed11dd26e71 100644 --- a/dpnp/backend/extensions/lapack/CMakeLists.txt +++ b/dpnp/backend/extensions/lapack/CMakeLists.txt @@ -69,7 +69,7 @@ if (DPNP_GENERATE_COVERAGE) target_link_options(${python_module_name} PRIVATE -fprofile-instr-generate -fcoverage-mapping) endif() -target_link_libraries(${python_module_name} PUBLIC MKL::MKL_DPCPP) +target_link_libraries(${python_module_name} PUBLIC MKL::MKL_SYCL::LAPACK) install(TARGETS ${python_module_name} DESTINATION "dpnp/backend/extensions/lapack" diff --git a/dpnp/backend/extensions/vm/CMakeLists.txt b/dpnp/backend/extensions/vm/CMakeLists.txt index a81dbe1df686..3df4d020f209 100644 --- a/dpnp/backend/extensions/vm/CMakeLists.txt +++ b/dpnp/backend/extensions/vm/CMakeLists.txt @@ -67,7 +67,7 @@ if (DPNP_GENERATE_COVERAGE) target_link_options(${python_module_name} PRIVATE -fprofile-instr-generate -fcoverage-mapping) endif() -target_link_libraries(${python_module_name} PUBLIC MKL::MKL_DPCPP) +target_link_libraries(${python_module_name} PUBLIC MKL::MKL_SYCL::VM) install(TARGETS ${python_module_name} DESTINATION "dpnp/backend/extensions/vm" diff --git a/dpnp/backend/extensions/vm/acos.hpp b/dpnp/backend/extensions/vm/acos.hpp new file mode 100644 index 000000000000..10bd68bf66e1 --- /dev/null +++ b/dpnp/backend/extensions/vm/acos.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event acos_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::acos(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AcosContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AcosOutputType::value_type, void>) + { + return nullptr; + } + else { + return acos_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/acosh.hpp b/dpnp/backend/extensions/vm/acosh.hpp new file mode 100644 index 000000000000..63d17cbd464b --- /dev/null +++ b/dpnp/backend/extensions/vm/acosh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event acosh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::acosh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AcoshContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AcoshOutputType::value_type, void>) + { + return nullptr; + } + else { + return acosh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/asin.hpp b/dpnp/backend/extensions/vm/asin.hpp new file mode 100644 index 000000000000..8df471f79c12 --- /dev/null +++ b/dpnp/backend/extensions/vm/asin.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event asin_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::asin(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AsinContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AsinOutputType::value_type, void>) + { + return nullptr; + } + else { + return asin_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/asinh.hpp b/dpnp/backend/extensions/vm/asinh.hpp new file mode 100644 index 000000000000..4e8e242257ac --- /dev/null +++ b/dpnp/backend/extensions/vm/asinh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event asinh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::asinh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AsinhContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AsinhOutputType::value_type, void>) + { + return nullptr; + } + else { + return asinh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/atan.hpp b/dpnp/backend/extensions/vm/atan.hpp new file mode 100644 index 000000000000..35967bb720a3 --- /dev/null +++ b/dpnp/backend/extensions/vm/atan.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event atan_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::atan(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AtanContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AtanOutputType::value_type, void>) + { + return nullptr; + } + else { + return atan_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/atan2.hpp b/dpnp/backend/extensions/vm/atan2.hpp new file mode 100644 index 000000000000..bf3d95a9c7db --- /dev/null +++ b/dpnp/backend/extensions/vm/atan2.hpp @@ -0,0 +1,81 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event atan2_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + const char *in_b, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + const T *b = reinterpret_cast(in_b); + T *y = reinterpret_cast(out_y); + + return mkl_vm::atan2(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing 1st input vector of size n + b, // pointer `b` containing 2nd input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct Atan2ContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::Atan2OutputType::value_type, void>) + { + return nullptr; + } + else { + return atan2_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/atanh.hpp b/dpnp/backend/extensions/vm/atanh.hpp new file mode 100644 index 000000000000..6f7ee1064e8c --- /dev/null +++ b/dpnp/backend/extensions/vm/atanh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event atanh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::atanh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AtanhContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AtanhOutputType::value_type, void>) + { + return nullptr; + } + else { + return atanh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/cosh.hpp b/dpnp/backend/extensions/vm/cosh.hpp new file mode 100644 index 000000000000..4a02bf361a50 --- /dev/null +++ b/dpnp/backend/extensions/vm/cosh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event cosh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::cosh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct CoshContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::CoshOutputType::value_type, void>) + { + return nullptr; + } + else { + return cosh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/pow.hpp b/dpnp/backend/extensions/vm/pow.hpp new file mode 100644 index 000000000000..5b7d13ca94f8 --- /dev/null +++ b/dpnp/backend/extensions/vm/pow.hpp @@ -0,0 +1,81 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event pow_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + const char *in_b, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + const T *b = reinterpret_cast(in_b); + T *y = reinterpret_cast(out_y); + + return mkl_vm::pow(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing 1st input vector of size n + b, // pointer `b` containing 2nd input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct PowContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::PowOutputType::value_type, void>) + { + return nullptr; + } + else { + return pow_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/round.hpp b/dpnp/backend/extensions/vm/round.hpp new file mode 100644 index 000000000000..0760db4d3fc4 --- /dev/null +++ b/dpnp/backend/extensions/vm/round.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event round_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::rint(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct RoundContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::RoundOutputType::value_type, void>) + { + return nullptr; + } + else { + return round_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/sinh.hpp b/dpnp/backend/extensions/vm/sinh.hpp new file mode 100644 index 000000000000..8b8114ad54ed --- /dev/null +++ b/dpnp/backend/extensions/vm/sinh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event sinh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::sinh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct SinhContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::SinhOutputType::value_type, void>) + { + return nullptr; + } + else { + return sinh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/tan.hpp b/dpnp/backend/extensions/vm/tan.hpp new file mode 100644 index 000000000000..4a9d0cc5e2e9 --- /dev/null +++ b/dpnp/backend/extensions/vm/tan.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event tan_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::tan(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct TanContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::TanOutputType::value_type, void>) + { + return nullptr; + } + else { + return tan_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/tanh.hpp b/dpnp/backend/extensions/vm/tanh.hpp new file mode 100644 index 000000000000..ed1be8f8d527 --- /dev/null +++ b/dpnp/backend/extensions/vm/tanh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +// +// 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. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event tanh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::tanh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct TanhContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::TanhOutputType::value_type, void>) + { + return nullptr; + } + else { + return tanh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/types_matrix.hpp b/dpnp/backend/extensions/vm/types_matrix.hpp index e710a5b46734..bbf74ed8c863 100644 --- a/dpnp/backend/extensions/vm/types_matrix.hpp +++ b/dpnp/backend/extensions/vm/types_matrix.hpp @@ -43,6 +43,40 @@ namespace vm { namespace types { +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::acos function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AcosOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::acosh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AcoshOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::add function. @@ -68,6 +102,89 @@ struct AddOutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::asin function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AsinOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::asinh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AsinhOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::atan function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AtanOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::atan2 function. + * + * @tparam T Type of input vectors `a` and `b` and of result vector `y`. + */ +template +struct Atan2OutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::BinaryTypeMapResultEntry, + dpctl_td_ns::BinaryTypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::atanh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AtanhOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::ceil function. @@ -110,10 +227,25 @@ template struct CosOutputType { using value_type = typename std::disjunction< - dpctl_td_ns:: - TypeMapResultEntry, std::complex>, - dpctl_td_ns:: - TypeMapResultEntry, std::complex>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::cosh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct CoshOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, dpctl_td_ns::TypeMapResultEntry, dpctl_td_ns::TypeMapResultEntry, dpctl_td_ns::DefaultResultEntry>::result_type; @@ -203,6 +335,46 @@ struct MulOutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::pow function. + * + * @tparam T Type of input vectors `a` and `b` and of result vector `y`. + */ +template +struct PowOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::BinaryTypeMapResultEntry, + T, + std::complex, + std::complex>, + dpctl_td_ns::BinaryTypeMapResultEntry, + T, + std::complex, + std::complex>, + dpctl_td_ns::BinaryTypeMapResultEntry, + dpctl_td_ns::BinaryTypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::rint function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct RoundOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::sin function. @@ -213,10 +385,25 @@ template struct SinOutputType { using value_type = typename std::disjunction< - dpctl_td_ns:: - TypeMapResultEntry, std::complex>, - dpctl_td_ns:: - TypeMapResultEntry, std::complex>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::sinh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct SinhOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, dpctl_td_ns::TypeMapResultEntry, dpctl_td_ns::TypeMapResultEntry, dpctl_td_ns::DefaultResultEntry>::result_type; @@ -281,6 +468,40 @@ struct SubOutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::tan function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct TanOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::tanh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct TanhOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::trunc function. diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index da088582af13..e4d89ff6e84f 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -30,19 +30,32 @@ #include #include +#include "acos.hpp" +#include "acosh.hpp" #include "add.hpp" +#include "asin.hpp" +#include "asinh.hpp" +#include "atan.hpp" +#include "atan2.hpp" +#include "atanh.hpp" #include "ceil.hpp" #include "common.hpp" #include "conj.hpp" #include "cos.hpp" +#include "cosh.hpp" #include "div.hpp" #include "floor.hpp" #include "ln.hpp" #include "mul.hpp" +#include "pow.hpp" +#include "round.hpp" #include "sin.hpp" +#include "sinh.hpp" #include "sqr.hpp" #include "sqrt.hpp" #include "sub.hpp" +#include "tan.hpp" +#include "tanh.hpp" #include "trunc.hpp" #include "types_matrix.hpp" @@ -52,18 +65,31 @@ namespace vm_ext = dpnp::backend::ext::vm; using vm_ext::binary_impl_fn_ptr_t; using vm_ext::unary_impl_fn_ptr_t; +static unary_impl_fn_ptr_t acos_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t acosh_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t add_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t asin_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t asinh_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t atan_dispatch_vector[dpctl_td_ns::num_types]; +static binary_impl_fn_ptr_t atan2_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t atanh_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t ceil_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t cos_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t cosh_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t div_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t floor_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t conj_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t ln_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t mul_dispatch_vector[dpctl_td_ns::num_types]; +static binary_impl_fn_ptr_t pow_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t round_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sin_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t sinh_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sqr_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sqrt_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t sub_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t tan_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t tanh_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t trunc_dispatch_vector[dpctl_td_ns::num_types]; PYBIND11_MODULE(_vm_impl, m) @@ -71,6 +97,62 @@ PYBIND11_MODULE(_vm_impl, m) using arrayT = dpctl::tensor::usm_ndarray; using event_vecT = std::vector; + // UnaryUfunc: ==== Acos(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + acos_dispatch_vector); + + auto acos_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + acos_dispatch_vector); + }; + m.def("_acos", acos_pyapi, + "Call `acos` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto acos_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + acos_dispatch_vector); + }; + m.def("_mkl_acos_to_call", acos_need_to_call_pyapi, + "Check input arguments to answer if `acos` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + + // UnaryUfunc: ==== Acosh(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + acosh_dispatch_vector); + + auto acosh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + acosh_dispatch_vector); + }; + m.def("_acosh", acosh_pyapi, + "Call `acosh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto acosh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + acosh_dispatch_vector); + }; + m.def("_mkl_acosh_to_call", acosh_need_to_call_pyapi, + "Check input arguments to answer if `acosh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // BinaryUfunc: ==== Add(x1, x2) ==== { vm_ext::init_ufunc_dispatch_vector( + asin_dispatch_vector); + + auto asin_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + asin_dispatch_vector); + }; + m.def("_asin", asin_pyapi, + "Call `asin` function from OneMKL VM library to compute " + "inverse sine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto asin_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + asin_dispatch_vector); + }; + m.def("_mkl_asin_to_call", asin_need_to_call_pyapi, + "Check input arguments to answer if `asin` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + + // UnaryUfunc: ==== Asinh(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + asinh_dispatch_vector); + + auto asinh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + asinh_dispatch_vector); + }; + m.def("_asinh", asinh_pyapi, + "Call `asinh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto asinh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + asinh_dispatch_vector); + }; + m.def("_mkl_asinh_to_call", asinh_need_to_call_pyapi, + "Check input arguments to answer if `asinh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + + // UnaryUfunc: ==== Atan(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + atan_dispatch_vector); + + auto atan_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + atan_dispatch_vector); + }; + m.def("_atan", atan_pyapi, + "Call `atan` function from OneMKL VM library to compute " + "inverse tangent of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto atan_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + atan_dispatch_vector); + }; + m.def("_mkl_atan_to_call", atan_need_to_call_pyapi, + "Check input arguments to answer if `atan` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + + // BinaryUfunc: ==== Atan2(x1, x2) ==== + { + vm_ext::init_ufunc_dispatch_vector( + atan2_dispatch_vector); + + auto atan2_pyapi = [&](sycl::queue exec_q, arrayT src1, arrayT src2, + arrayT dst, const event_vecT &depends = {}) { + return vm_ext::binary_ufunc(exec_q, src1, src2, dst, depends, + atan2_dispatch_vector); + }; + m.def("_atan2", atan2_pyapi, + "Call `atan2` function from OneMKL VM library to compute element " + "by element inverse tangent of `x1/x2`", + py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), + py::arg("dst"), py::arg("depends") = py::list()); + + auto atan2_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src1, + arrayT src2, arrayT dst) { + return vm_ext::need_to_call_binary_ufunc(exec_q, src1, src2, dst, + atan2_dispatch_vector); + }; + m.def("_mkl_atan2_to_call", atan2_need_to_call_pyapi, + "Check input arguments to answer if `atan2` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), + py::arg("dst")); + } + + // UnaryUfunc: ==== Atanh(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + atanh_dispatch_vector); + + auto atanh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + atanh_dispatch_vector); + }; + m.def("_atanh", atanh_pyapi, + "Call `atanh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto atanh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + atanh_dispatch_vector); + }; + m.def("_mkl_atanh_to_call", atanh_need_to_call_pyapi, + "Check input arguments to answer if `atanh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Ceil(x) ==== { vm_ext::init_ufunc_dispatch_vector( + cosh_dispatch_vector); + + auto cosh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + cosh_dispatch_vector); + }; + m.def("_cosh", cosh_pyapi, + "Call `cosh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto cosh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + cosh_dispatch_vector); + }; + m.def("_mkl_cosh_to_call", cosh_need_to_call_pyapi, + "Check input arguments to answer if `cosh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // BinaryUfunc: ==== Div(x1, x2) ==== { vm_ext::init_ufunc_dispatch_vector( + pow_dispatch_vector); + + auto pow_pyapi = [&](sycl::queue exec_q, arrayT src1, arrayT src2, + arrayT dst, const event_vecT &depends = {}) { + return vm_ext::binary_ufunc(exec_q, src1, src2, dst, depends, + pow_dispatch_vector); + }; + m.def("_pow", pow_pyapi, + "Call `pow` function from OneMKL VM library to performs element " + "by element exponentiation of vector `src1` raised to the power " + "of vector `src2` to resulting vector `dst`", + py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), + py::arg("dst"), py::arg("depends") = py::list()); + + auto pow_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src1, + arrayT src2, arrayT dst) { + return vm_ext::need_to_call_binary_ufunc(exec_q, src1, src2, dst, + pow_dispatch_vector); + }; + m.def("_mkl_pow_to_call", pow_need_to_call_pyapi, + "Check input arguments to answer if `pow` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), + py::arg("dst")); + } + + // UnaryUfunc: ==== Round(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + round_dispatch_vector); + + auto round_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + round_dispatch_vector); + }; + m.def("_round", round_pyapi, + "Call `rint` function from OneMKL VM library to compute " + "the rounded value of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto round_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + round_dispatch_vector); + }; + m.def("_mkl_round_to_call", round_need_to_call_pyapi, + "Check input arguments to answer if `rint` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Sin(x) ==== { vm_ext::init_ufunc_dispatch_vector( + sinh_dispatch_vector); + + auto sinh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + sinh_dispatch_vector); + }; + m.def("_sinh", sinh_pyapi, + "Call `sinh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto sinh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + sinh_dispatch_vector); + }; + m.def("_mkl_sinh_to_call", sinh_need_to_call_pyapi, + "Check input arguments to answer if `sinh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Sqr(x) ==== { vm_ext::init_ufunc_dispatch_vector( + tan_dispatch_vector); + + auto tan_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + tan_dispatch_vector); + }; + m.def("_tan", tan_pyapi, + "Call `tan` function from OneMKL VM library to compute " + "tangent of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto tan_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + tan_dispatch_vector); + }; + m.def("_mkl_tan_to_call", tan_need_to_call_pyapi, + "Check input arguments to answer if `tan` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + + // UnaryUfunc: ==== Tanh(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + tanh_dispatch_vector); + + auto tanh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + tanh_dispatch_vector); + }; + m.def("_tanh", tanh_pyapi, + "Call `tanh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto tanh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + tanh_dispatch_vector); + }; + m.def("_mkl_tanh_to_call", tanh_need_to_call_pyapi, + "Check input arguments to answer if `tanh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Trunc(x) ==== { vm_ext::init_ufunc_dispatch_vector(vec_sz) * max_sg_size < size) { - using multi_ptrT = - sycl::multi_ptr<_DataType, - sycl::access::address_space::global_space>; + auto input_multi_ptr = sycl::address_space_cast< + sycl::access::address_space::global_space, + sycl::access::decorated::yes>(&input_data[start]); + auto result_multi_ptr = sycl::address_space_cast< + sycl::access::address_space::global_space, + sycl::access::decorated::yes>(&result[start]); - sycl::vec<_DataType, vec_sz> x = - sg.load(multi_ptrT(&input_data[start])); + sycl::vec<_DataType, vec_sz> x = sg.load(input_multi_ptr); sycl::vec<_DataType, vec_sz> res_vec; if constexpr (std::is_same_v<_DataType, bool>) { @@ -86,7 +88,7 @@ DPCTLSyclEventRef dpnp_invert_c(DPCTLSyclQueueRef q_ref, res_vec = ~x; } - sg.store(multi_ptrT(&result[start]), res_vec); + sg.store(result_multi_ptr, res_vec); } else { for (size_t k = start + sg.get_local_id()[0]; k < size; diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 8fb510e9fab1..4b26da44853e 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -233,15 +233,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCCOS][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_acos_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_acos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_acos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_acos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_acos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_acosh_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCCOSH][eft_LNG][eft_LNG] = { @@ -251,15 +242,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCCOSH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_acosh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_acosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_acosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_acosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_acosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_asin_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCSIN][eft_LNG][eft_LNG] = { @@ -269,15 +251,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCSIN][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_asin_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_asin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_asin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_asin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_asin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_asinh_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCSINH][eft_LNG][eft_LNG] = { @@ -287,15 +260,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCSINH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_asinh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_asinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_asinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_asinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_asinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_atan_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCTAN][eft_LNG][eft_LNG] = { @@ -305,15 +269,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCTAN][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_atan_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_atan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_atan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_atan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_atan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_atanh_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCTANH][eft_LNG][eft_LNG] = { @@ -323,15 +278,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCTANH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_atanh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_atanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_atanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_atanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_atanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_CBRT][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_cbrt_c_default}; fmap[DPNPFuncName::DPNP_FN_CBRT][eft_LNG][eft_LNG] = { @@ -453,15 +399,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_COSH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_cosh_c_default}; - fmap[DPNPFuncName::DPNP_FN_COSH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_cosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COSH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_cosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COSH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_cosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COSH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_degrees_c_default}; fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_LNG][eft_LNG] = { @@ -660,15 +597,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_SINH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_sinh_c_default}; - fmap[DPNPFuncName::DPNP_FN_SINH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_sinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SINH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_sinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SINH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_sinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SINH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_sinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SQRT][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_sqrt_c_default}; fmap[DPNPFuncName::DPNP_FN_SQRT][eft_LNG][eft_LNG] = { @@ -693,15 +621,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_TAN][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_tan_c_default}; - fmap[DPNPFuncName::DPNP_FN_TAN_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_tan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TAN_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_tan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TAN_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_tan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TAN_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_tan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TANH][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_tanh_c_default}; fmap[DPNPFuncName::DPNP_FN_TANH][eft_LNG][eft_LNG] = { @@ -711,15 +630,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_TANH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_tanh_c_default}; - fmap[DPNPFuncName::DPNP_FN_TANH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_tanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TANH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_tanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TANH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_tanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TANH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_tanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_trunc_c_default}; fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_LNG][eft_LNG] = { @@ -1326,8 +1236,6 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) { \ constexpr size_t lws = 64; \ constexpr unsigned int vec_sz = 8; \ - constexpr sycl::access::address_space global_space = \ - sycl::access::address_space::global_space; \ \ auto gws_range = sycl::range<1>( \ ((result_size + lws * vec_sz - 1) / (lws * vec_sz)) * \ @@ -1344,12 +1252,17 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) \ if (start + static_cast(vec_sz) * max_sg_size < \ result_size) { \ - using input1_ptrT = \ - sycl::multi_ptr<_DataType_input1, global_space>; \ - using input2_ptrT = \ - sycl::multi_ptr<_DataType_input2, global_space>; \ - using result_ptrT = \ - sycl::multi_ptr<_DataType_output, global_space>; \ + auto input1_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>( \ + &input1_data[start]); \ + auto input2_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>( \ + &input2_data[start]); \ + auto result_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>(&result[start]); \ \ sycl::vec<_DataType_output, vec_sz> res_vec; \ \ @@ -1363,11 +1276,9 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) _DataType_output>) \ { \ sycl::vec<_DataType_input1, vec_sz> x1 = \ - sg.load( \ - input1_ptrT(&input1_data[start])); \ + sg.load(input1_multi_ptr); \ sycl::vec<_DataType_input2, vec_sz> x2 = \ - sg.load( \ - input2_ptrT(&input2_data[start])); \ + sg.load(input2_multi_ptr); \ \ res_vec = __vec_operation__; \ } \ @@ -1377,24 +1288,20 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) sycl::vec<_DataType_output, vec_sz> x1 = \ dpnp_vec_cast<_DataType_output, \ _DataType_input1, vec_sz>( \ - sg.load(input1_ptrT( \ - &input1_data[start]))); \ + sg.load(input1_multi_ptr)); \ sycl::vec<_DataType_output, vec_sz> x2 = \ dpnp_vec_cast<_DataType_output, \ _DataType_input2, vec_sz>( \ - sg.load(input2_ptrT( \ - &input2_data[start]))); \ + sg.load(input2_multi_ptr)); \ \ res_vec = __vec_operation__; \ } \ } \ else { \ sycl::vec<_DataType_input1, vec_sz> x1 = \ - sg.load( \ - input1_ptrT(&input1_data[start])); \ + sg.load(input1_multi_ptr); \ sycl::vec<_DataType_input2, vec_sz> x2 = \ - sg.load( \ - input2_ptrT(&input2_data[start])); \ + sg.load(input2_multi_ptr); \ \ for (size_t k = 0; k < vec_sz; ++k) { \ const _DataType_output input1_elem = x1[k]; \ @@ -1402,8 +1309,7 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) res_vec[k] = __operation__; \ } \ } \ - sg.store(result_ptrT(&result[start]), \ - res_vec); \ + sg.store(result_multi_ptr, res_vec); \ } \ else { \ for (size_t k = start + sg.get_local_id()[0]; \ @@ -1569,13 +1475,6 @@ static void func_map_elemwise_2arg_3type_core(func_map_t &fmap) func_type_map_t::find_type, func_type_map_t::find_type>}), ...); - ((fmap[DPNPFuncName::DPNP_FN_POWER_EXT][FT1][FTs] = - {populate_func_types(), - (void *)dpnp_power_c_ext< - func_type_map_t::find_type()>, - func_type_map_t::find_type, - func_type_map_t::find_type>}), - ...); ((fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][FT1][FTs] = {populate_func_types(), (void *)dpnp_subtract_c_ext< @@ -1588,19 +1487,6 @@ static void func_map_elemwise_2arg_3type_core(func_map_t &fmap) template static void func_map_elemwise_2arg_3type_short_core(func_map_t &fmap) { - ((fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][FT1][FTs] = - {get_floating_res_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type()>, - func_type_map_t::find_type, - func_type_map_t::find_type>, - get_floating_res_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type< - get_floating_res_type()>, - func_type_map_t::find_type, - func_type_map_t::find_type>}), - ...); ((fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][FT1][FTs] = {get_floating_res_type(), (void *)dpnp_copysign_c_ext< diff --git a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp index b8fa2179b6f6..2cca84f9e61f 100644 --- a/dpnp/backend/kernels/dpnp_krnl_indexing.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_indexing.cpp @@ -602,17 +602,6 @@ void (*dpnp_put_default_c)(void *, const size_t) = dpnp_put_c<_DataType, _IndecesType, _ValueType>; -template -DPCTLSyclEventRef (*dpnp_put_ext_c)(DPCTLSyclQueueRef, - void *, - void *, - void *, - const size_t, - const size_t, - const size_t, - const DPCTLEventVectorRef) = - dpnp_put_c<_DataType, _IndecesType, _ValueType>; - template DPCTLSyclEventRef dpnp_put_along_axis_c(DPCTLSyclQueueRef q_ref, @@ -1007,15 +996,6 @@ void func_map_init_indexing_func(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_PUT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_put_default_c}; - fmap[DPNPFuncName::DPNP_FN_PUT_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_put_ext_c}; - fmap[DPNPFuncName::DPNP_FN_PUT_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_put_ext_c}; - fmap[DPNPFuncName::DPNP_FN_PUT_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_put_ext_c}; - fmap[DPNPFuncName::DPNP_FN_PUT_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_put_ext_c}; - fmap[DPNPFuncName::DPNP_FN_PUT_ALONG_AXIS][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_put_along_axis_default_c}; fmap[DPNPFuncName::DPNP_FN_PUT_ALONG_AXIS][eft_LNG][eft_LNG] = { diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index 1757d053416a..2080306aeade 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -74,7 +74,7 @@ DPCTLSyclEventRef dpnp_all_c(DPCTLSyclQueueRef q_ref, sycl::nd_range<1> gws(gws_range, lws_range); auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { - auto gr = nd_it.get_group(); + auto gr = nd_it.get_sub_group(); const auto max_gr_size = gr.get_max_local_range()[0]; const size_t start = vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + @@ -127,8 +127,79 @@ DPCTLSyclEventRef (*dpnp_all_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_all_c<_DataType, _ResultType>; -template -class dpnp_allclose_c_kernel; +template +class dpnp_allclose_kernel; + +template +static sycl::event dpnp_allclose(sycl::queue &q, + const _DataType1 *array1, + const _DataType2 *array2, + bool *result, + const size_t size, + const _TolType rtol_val, + const _TolType atol_val) +{ + sycl::event fill_event = q.fill(result, true, 1); + if (!size) { + return fill_event; + } + + constexpr size_t lws = 64; + constexpr size_t vec_sz = 8; + + auto gws_range = + sycl::range<1>(((size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); + auto lws_range = sycl::range<1>(lws); + sycl::nd_range<1> gws(gws_range, lws_range); + + auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { + auto gr = nd_it.get_sub_group(); + const auto max_gr_size = gr.get_max_local_range()[0]; + const auto gr_size = gr.get_local_linear_range(); + const size_t start = + vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + + gr.get_group_linear_id() * max_gr_size); + const size_t end = sycl::min(start + vec_sz * gr_size, size); + + // each work-item iterates over "vec_sz" elements in the input arrays + bool partial = true; + + for (size_t i = start + gr.get_local_linear_id(); i < end; i += gr_size) + { + if constexpr (std::is_floating_point_v<_DataType1> && + std::is_floating_point_v<_DataType2>) + { + if (std::isinf(array1[i]) || std::isinf(array2[i])) { + partial &= (array1[i] == array2[i]); + continue; + } + } + + // casting integeral to floating type to avoid bad behavior + // on abs(MIN_INT), which leads to undefined result + using _Arr2Type = std::conditional_t, + _TolType, _DataType2>; + _Arr2Type arr2 = static_cast<_Arr2Type>(array2[i]); + + partial &= (std::abs(array1[i] - arr2) <= + (atol_val + rtol_val * std::abs(arr2))); + } + partial = sycl::all_of_group(gr, partial); + + if (gr.leader() && (partial == false)) { + result[0] = false; + } + }; + + auto kernel_func = [&](sycl::handler &cgh) { + cgh.depends_on(fill_event); + cgh.parallel_for< + class dpnp_allclose_kernel<_DataType1, _DataType2, _TolType>>( + gws, kernel_parallel_for_func); + }; + + return q.submit(kernel_func); +} template DPCTLSyclEventRef dpnp_allclose_c(DPCTLSyclQueueRef q_ref, @@ -140,6 +211,9 @@ DPCTLSyclEventRef dpnp_allclose_c(DPCTLSyclQueueRef q_ref, double atol_val, const DPCTLEventVectorRef dep_event_vec_ref) { + static_assert(std::is_same_v<_ResultType, bool>, + "Boolean result type is required"); + // avoid warning unused variable (void)dep_event_vec_ref; @@ -152,40 +226,21 @@ DPCTLSyclEventRef dpnp_allclose_c(DPCTLSyclQueueRef q_ref, sycl::queue q = *(reinterpret_cast(q_ref)); sycl::event event; - DPNPC_ptr_adapter<_DataType1> input1_ptr(q_ref, array1_in, size); - DPNPC_ptr_adapter<_DataType2> input2_ptr(q_ref, array2_in, size); - DPNPC_ptr_adapter<_ResultType> result1_ptr(q_ref, result1, 1, true, true); - const _DataType1 *array1 = input1_ptr.get_ptr(); - const _DataType2 *array2 = input2_ptr.get_ptr(); - _ResultType *result = result1_ptr.get_ptr(); - - result[0] = true; + const _DataType1 *array1 = static_cast(array1_in); + const _DataType2 *array2 = static_cast(array2_in); + bool *result = static_cast(result1); - if (!size) { - return event_ref; + if (q.get_device().has(sycl::aspect::fp64)) { + event = + dpnp_allclose(q, array1, array2, result, size, rtol_val, atol_val); + } + else { + float rtol = static_cast(rtol_val); + float atol = static_cast(atol_val); + event = dpnp_allclose(q, array1, array2, result, size, rtol, atol); } - - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - - if (std::abs(array1[i] - array2[i]) > - (atol_val + rtol_val * std::abs(array2[i]))) - { - result[0] = false; - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for< - class dpnp_allclose_c_kernel<_DataType1, _DataType2, _ResultType>>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); } @@ -269,7 +324,7 @@ DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref, sycl::nd_range<1> gws(gws_range, lws_range); auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { - auto gr = nd_it.get_group(); + auto gr = nd_it.get_sub_group(); const auto max_gr_size = gr.get_max_local_range()[0]; const size_t start = vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + @@ -521,8 +576,6 @@ DPCTLSyclEventRef (*dpnp_any_ext_c)(DPCTLSyclQueueRef, else { \ constexpr size_t lws = 64; \ constexpr unsigned int vec_sz = 8; \ - constexpr sycl::access::address_space global_space = \ - sycl::access::address_space::global_space; \ \ auto gws_range = sycl::range<1>( \ ((result_size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); \ @@ -537,12 +590,20 @@ DPCTLSyclEventRef (*dpnp_any_ext_c)(DPCTLSyclQueueRef, \ if (start + static_cast(vec_sz) * max_sg_size < \ result_size) { \ - sycl::vec<_DataType_input1, vec_sz> x1 = sg.load( \ - sycl::multi_ptr<_DataType_input1, global_space>( \ - &input1_data[start])); \ - sycl::vec<_DataType_input2, vec_sz> x2 = sg.load( \ - sycl::multi_ptr<_DataType_input2, global_space>( \ - &input2_data[start])); \ + auto input1_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>(&input1_data[start]); \ + auto input2_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>(&input2_data[start]); \ + auto result_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>(&result[start]); \ + \ + sycl::vec<_DataType_input1, vec_sz> x1 = \ + sg.load(input1_multi_ptr); \ + sycl::vec<_DataType_input2, vec_sz> x2 = \ + sg.load(input2_multi_ptr); \ sycl::vec res_vec; \ \ for (size_t k = 0; k < vec_sz; ++k) { \ @@ -550,9 +611,7 @@ DPCTLSyclEventRef (*dpnp_any_ext_c)(DPCTLSyclQueueRef, const _DataType_input2 input2_elem = x2[k]; \ res_vec[k] = __operation__; \ } \ - sg.store( \ - sycl::multi_ptr(&result[start]), \ - res_vec); \ + sg.store(result_multi_ptr, res_vec); \ } \ else { \ for (size_t k = start; k < result_size; ++k) { \ diff --git a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp index 4e29874b85f8..b4f5cd96b4d0 100644 --- a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp @@ -109,15 +109,6 @@ template void (*dpnp_around_default_c)(const void *, void *, const size_t, const int) = dpnp_around_c<_DataType>; -template -DPCTLSyclEventRef (*dpnp_around_ext_c)(DPCTLSyclQueueRef, - const void *, - void *, - const size_t, - const int, - const DPCTLEventVectorRef) = - dpnp_around_c<_DataType>; - template class dpnp_elemwise_absolute_c_kernel; @@ -160,8 +151,6 @@ DPCTLSyclEventRef constexpr size_t lws = 64; constexpr unsigned int vec_sz = 8; - constexpr sycl::access::address_space global_space = - sycl::access::address_space::global_space; auto gws_range = sycl::range<1>(((size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); @@ -175,18 +164,20 @@ DPCTLSyclEventRef sg.get_group_id()[0] * max_sg_size); if (start + static_cast(vec_sz) * max_sg_size < size) { - using input_ptrT = - sycl::multi_ptr<_DataType_input, global_space>; - using result_ptrT = - sycl::multi_ptr<_DataType_output, global_space>; + auto array_multi_ptr = sycl::address_space_cast< + sycl::access::address_space::global_space, + sycl::access::decorated::yes>(&array1[start]); + auto result_multi_ptr = sycl::address_space_cast< + sycl::access::address_space::global_space, + sycl::access::decorated::yes>(&result[start]); sycl::vec<_DataType_input, vec_sz> data_vec = - sg.load(input_ptrT(&array1[start])); + sg.load(array_multi_ptr); sycl::vec<_DataType_output, vec_sz> res_vec = sycl::abs(data_vec); - sg.store(result_ptrT(&result[start]), res_vec); + sg.store(result_multi_ptr, res_vec); } else { for (size_t k = start + sg.get_local_id()[0]; k < size; @@ -1184,15 +1175,6 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_AROUND][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_around_default_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_around_ext_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_around_ext_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_around_ext_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_around_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_cross_default_c}; fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_LNG] = { diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 9a4a556435ac..a512737ef31e 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -38,28 +38,12 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_ALLCLOSE DPNP_FN_ALLCLOSE_EXT DPNP_FN_ARANGE - DPNP_FN_ARCCOS - DPNP_FN_ARCCOS_EXT - DPNP_FN_ARCCOSH - DPNP_FN_ARCCOSH_EXT - DPNP_FN_ARCSIN - DPNP_FN_ARCSIN_EXT - DPNP_FN_ARCSINH - DPNP_FN_ARCSINH_EXT - DPNP_FN_ARCTAN - DPNP_FN_ARCTAN_EXT - DPNP_FN_ARCTAN2 - DPNP_FN_ARCTAN2_EXT - DPNP_FN_ARCTANH - DPNP_FN_ARCTANH_EXT DPNP_FN_ARGMAX DPNP_FN_ARGMAX_EXT DPNP_FN_ARGMIN DPNP_FN_ARGMIN_EXT DPNP_FN_ARGSORT DPNP_FN_ARGSORT_EXT - DPNP_FN_AROUND - DPNP_FN_AROUND_EXT DPNP_FN_ASTYPE DPNP_FN_ASTYPE_EXT DPNP_FN_CBRT @@ -74,8 +58,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_COPYSIGN_EXT DPNP_FN_CORRELATE DPNP_FN_CORRELATE_EXT - DPNP_FN_COSH - DPNP_FN_COSH_EXT DPNP_FN_COUNT_NONZERO DPNP_FN_COUNT_NONZERO_EXT DPNP_FN_CROSS @@ -130,8 +112,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_HYPOT_EXT DPNP_FN_IDENTITY DPNP_FN_IDENTITY_EXT - DPNP_FN_INITVAL - DPNP_FN_INITVAL_EXT DPNP_FN_INV DPNP_FN_INV_EXT DPNP_FN_KRON @@ -166,14 +146,10 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_PARTITION DPNP_FN_PARTITION_EXT DPNP_FN_PLACE - DPNP_FN_POWER - DPNP_FN_POWER_EXT DPNP_FN_PROD DPNP_FN_PROD_EXT DPNP_FN_PTP DPNP_FN_PTP_EXT - DPNP_FN_PUT - DPNP_FN_PUT_EXT DPNP_FN_QR DPNP_FN_QR_EXT DPNP_FN_RADIANS @@ -255,8 +231,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_RNG_ZIPF_EXT DPNP_FN_SEARCHSORTED DPNP_FN_SEARCHSORTED_EXT - DPNP_FN_SINH - DPNP_FN_SINH_EXT DPNP_FN_SORT DPNP_FN_SORT_EXT DPNP_FN_STD @@ -265,10 +239,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_SUM_EXT DPNP_FN_SVD DPNP_FN_SVD_EXT - DPNP_FN_TAN - DPNP_FN_TAN_EXT - DPNP_FN_TANH - DPNP_FN_TANH_EXT DPNP_FN_TRACE DPNP_FN_TRACE_EXT DPNP_FN_TRANSPOSE @@ -409,22 +379,17 @@ cpdef dpnp_descriptor dpnp_matmul(dpnp_descriptor in_array1, dpnp_descriptor in_ """ Array creation routines """ -cpdef dpnp_descriptor dpnp_init_val(shape, dtype, value) cpdef dpnp_descriptor dpnp_copy(dpnp_descriptor x1) """ Mathematical functions """ -cpdef dpnp_descriptor dpnp_arctan2(dpnp_descriptor x1_obj, dpnp_descriptor x2_obj, object dtype=*, - dpnp_descriptor out=*, object where=*) cpdef dpnp_descriptor dpnp_hypot(dpnp_descriptor x1_obj, dpnp_descriptor x2_obj, object dtype=*, dpnp_descriptor out=*, object where=*) cpdef dpnp_descriptor dpnp_maximum(dpnp_descriptor x1_obj, dpnp_descriptor x2_obj, object dtype=*, dpnp_descriptor out=*, object where=*) cpdef dpnp_descriptor dpnp_minimum(dpnp_descriptor x1_obj, dpnp_descriptor x2_obj, object dtype=*, dpnp_descriptor out=*, object where=*) -cpdef dpnp_descriptor dpnp_power(dpnp_descriptor x1_obj, dpnp_descriptor x2_obj, object dtype=*, - dpnp_descriptor out=*, object where=*) """ Array manipulation routines @@ -453,14 +418,7 @@ cpdef dpnp_descriptor dpnp_argmin(dpnp_descriptor array1) """ Trigonometric functions """ -cpdef dpnp_descriptor dpnp_arccos(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_arccosh(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_arcsin(dpnp_descriptor array1, dpnp_descriptor out) -cpdef dpnp_descriptor dpnp_arcsinh(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_arctan(dpnp_descriptor array1, dpnp_descriptor out) -cpdef dpnp_descriptor dpnp_arctanh(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_cbrt(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_cosh(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_degrees(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_exp(dpnp_descriptor array1, dpnp_descriptor out) cpdef dpnp_descriptor dpnp_exp2(dpnp_descriptor array1) @@ -470,6 +428,3 @@ cpdef dpnp_descriptor dpnp_log1p(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_log2(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_radians(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_recip(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_sinh(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_tan(dpnp_descriptor array1, dpnp_descriptor out) -cpdef dpnp_descriptor dpnp_tanh(dpnp_descriptor array1) diff --git a/dpnp/dpnp_algo/dpnp_algo.pyx b/dpnp/dpnp_algo/dpnp_algo.pyx index c697e76c682b..0bde1b25fb2e 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pyx +++ b/dpnp/dpnp_algo/dpnp_algo.pyx @@ -56,7 +56,6 @@ import numpy __all__ = [ "dpnp_astype", "dpnp_flatten", - "dpnp_init_val", "dpnp_queue_initialize", ] @@ -85,9 +84,6 @@ ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_flatten_t)(c_dpctl.DPCTLSyclQueueR const shape_elem_type * , const shape_elem_type * , const long * , const c_dpctl.DPCTLEventVectorRef) -ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_initval_t)(c_dpctl.DPCTLSyclQueueRef, - void *, void * , size_t, - const c_dpctl.DPCTLEventVectorRef) cpdef utils.dpnp_descriptor dpnp_astype(utils.dpnp_descriptor x1, dtype): @@ -168,39 +164,6 @@ cpdef utils.dpnp_descriptor dpnp_flatten(utils.dpnp_descriptor x1): return result -cpdef utils.dpnp_descriptor dpnp_init_val(shape, dtype, value): - """ - same as dpnp_full(). TODO remove code dumplication - """ - cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype) - - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_INITVAL_EXT, param1_type, param1_type) - - cdef utils.dpnp_descriptor result = utils_py.create_output_descriptor_py(shape, dtype, None) - - result_obj = result.get_array() - - # TODO: find better way to pass single value with type conversion - cdef utils.dpnp_descriptor val_arr = utils_py.create_output_descriptor_py((1, ), - dtype, - None, - device=result_obj.sycl_device, - usm_type=result_obj.usm_type, - sycl_queue=result_obj.sycl_queue) - val_arr.get_pyobj()[0] = value - - cdef c_dpctl.SyclQueue q = result_obj.sycl_queue - cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() - - cdef fptr_dpnp_initval_t func = kernel_data.ptr - cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), val_arr.get_data(), result.size, NULL) - - with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) - c_dpctl.DPCTLEvent_Delete(event_ref) - - return result - - cpdef dpnp_queue_initialize(): """ Initialize SYCL queue which will be used for any library operations. @@ -219,7 +182,7 @@ cpdef dpnp_queue_initialize(): # TODO: # choose seed number as is in numpy seed_from_time = time(NULL) - dpnp_rng_srand_c(seed_from_time) + dpnp_rng_srand_c(< size_t > seed_from_time) """ diff --git a/dpnp/dpnp_algo/dpnp_algo_indexing.pxi b/dpnp/dpnp_algo/dpnp_algo_indexing.pxi index 808961298c22..36fc7ff8eb91 100644 --- a/dpnp/dpnp_algo/dpnp_algo_indexing.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_indexing.pxi @@ -41,7 +41,6 @@ __all__ += [ "dpnp_diagonal", "dpnp_fill_diagonal", "dpnp_indices", - "dpnp_put", "dpnp_put_along_axis", "dpnp_putmask", "dpnp_select", @@ -80,14 +79,6 @@ ctypedef c_dpctl.DPCTLSyclEventRef(*custom_indexing_3in_with_axis_func_ptr_t)(c_ const size_t, const size_t, const c_dpctl.DPCTLEventVectorRef) -ctypedef c_dpctl.DPCTLSyclEventRef(*custom_indexing_6in_func_ptr_t)(c_dpctl.DPCTLSyclQueueRef, - void *, - void * , - void * , - const size_t, - const size_t, - const size_t, - const c_dpctl.DPCTLEventVectorRef) cpdef utils.dpnp_descriptor dpnp_choose(utils.dpnp_descriptor x1, list choices1): @@ -292,65 +283,6 @@ cpdef object dpnp_indices(dimensions): return dpnp_result -cpdef dpnp_put(dpnp_descriptor x1, object ind, v): - ind_is_list = isinstance(ind, list) - - x1_obj = x1.get_array() - - if dpnp.isscalar(ind): - ind_size = 1 - else: - ind_size = len(ind) - cdef utils.dpnp_descriptor ind_array = utils_py.create_output_descriptor_py((ind_size,), - dpnp.int64, - None, - device=x1_obj.sycl_device, - usm_type=x1_obj.usm_type, - sycl_queue=x1_obj.sycl_queue) - if dpnp.isscalar(ind): - ind_array.get_pyobj()[0] = ind - else: - for i in range(ind_size): - ind_array.get_pyobj()[i] = ind[i] - - if dpnp.isscalar(v): - v_size = 1 - else: - v_size = len(v) - cdef utils.dpnp_descriptor v_array = utils_py.create_output_descriptor_py((v_size,), - x1.dtype, - None, - device=x1_obj.sycl_device, - usm_type=x1_obj.usm_type, - sycl_queue=x1_obj.sycl_queue) - if dpnp.isscalar(v): - v_array.get_pyobj()[0] = v - else: - for i in range(v_size): - v_array.get_pyobj()[i] = v[i] - - cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype) - - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_PUT_EXT, param1_type, param1_type) - - cdef c_dpctl.SyclQueue q = x1_obj.sycl_queue - cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() - - cdef custom_indexing_6in_func_ptr_t func = kernel_data.ptr - - cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, - x1.get_data(), - ind_array.get_data(), - v_array.get_data(), - x1.size, - ind_array.size, - v_array.size, - NULL) # dep_events_ref - - with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) - c_dpctl.DPCTLEvent_Delete(event_ref) - - cpdef dpnp_put_along_axis(dpnp_descriptor arr, dpnp_descriptor indices, dpnp_descriptor values, int axis): cdef shape_type_c arr_shape = arr.shape cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(arr.dtype) diff --git a/dpnp/dpnp_algo/dpnp_algo_manipulation.pxi b/dpnp/dpnp_algo/dpnp_algo_manipulation.pxi index 43b8310578df..4280bd621ce9 100644 --- a/dpnp/dpnp_algo/dpnp_algo_manipulation.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_manipulation.pxi @@ -38,9 +38,7 @@ and the rest of the library __all__ += [ "dpnp_atleast_2d", "dpnp_atleast_3d", - "dpnp_expand_dims", "dpnp_repeat", - "dpnp_reshape", ] @@ -104,35 +102,6 @@ cpdef utils.dpnp_descriptor dpnp_atleast_3d(utils.dpnp_descriptor arr): return arr -cpdef utils.dpnp_descriptor dpnp_expand_dims(utils.dpnp_descriptor in_array, axis): - axis_tuple = utils._object_to_tuple(axis) - result_ndim = len(axis_tuple) + in_array.ndim - - if len(axis_tuple) == 0: - axis_ndim = 0 - else: - axis_ndim = max(-min(0, min(axis_tuple)), max(0, max(axis_tuple))) + 1 - - axis_norm = utils._object_to_tuple(utils.normalize_axis(axis_tuple, result_ndim)) - - if axis_ndim - len(axis_norm) > in_array.ndim: - utils.checker_throw_axis_error("dpnp_expand_dims", "axis", axis, axis_ndim) - - if len(axis_norm) > len(set(axis_norm)): - utils.checker_throw_value_error("dpnp_expand_dims", "axis", axis, "no repeated axis") - - cdef shape_type_c shape_list - axis_idx = 0 - for i in range(result_ndim): - if i in axis_norm: - shape_list.push_back(1) - else: - shape_list.push_back(in_array.shape[axis_idx]) - axis_idx = axis_idx + 1 - - return dpnp_reshape(in_array, shape_list) - - cpdef utils.dpnp_descriptor dpnp_repeat(utils.dpnp_descriptor array1, repeats, axes=None): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(array1.dtype) @@ -165,17 +134,3 @@ cpdef utils.dpnp_descriptor dpnp_repeat(utils.dpnp_descriptor array1, repeats, a c_dpctl.DPCTLEvent_Delete(event_ref) return result - - -cpdef utils.dpnp_descriptor dpnp_reshape(utils.dpnp_descriptor array1, newshape, order="C"): - # return dpnp.get_dpnp_descriptor(dpctl.tensor.usm_ndarray(newshape, dtype=numpy.dtype(array1.dtype).name, buffer=array1.get_pyobj())) - # return dpnp.get_dpnp_descriptor(dpctl.tensor.reshape(array1.get_pyobj(), newshape)) - array1_obj = array1.get_array() - array_obj = dpctl.tensor.reshape(array1_obj, newshape, order=order) - return dpnp.get_dpnp_descriptor(dpnp_array(array_obj.shape, - buffer=array_obj, - order=order, - device=array1_obj.sycl_device, - usm_type=array1_obj.usm_type, - sycl_queue=array1_obj.sycl_queue), - copy_when_nondefault_queue=False) diff --git a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi index 86579933b018..401232ca8f25 100644 --- a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi @@ -37,8 +37,6 @@ and the rest of the library __all__ += [ "dpnp_absolute", - "dpnp_arctan2", - "dpnp_around", "dpnp_copysign", "dpnp_cross", "dpnp_cumprod", @@ -56,7 +54,6 @@ __all__ += [ "dpnp_nancumsum", "dpnp_nanprod", "dpnp_nansum", - "dpnp_power", "dpnp_prod", "dpnp_sum", "dpnp_trapz", @@ -72,9 +69,6 @@ ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_1in_2out_t)(c_dpctl.DPCTLSyclQueueRef, ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_trapz_2in_1out_with_2size_t)(c_dpctl.DPCTLSyclQueueRef, void *, void * , void * , double, size_t, size_t, const c_dpctl.DPCTLEventVectorRef) -ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_around_1in_1out_t)(c_dpctl.DPCTLSyclQueueRef, - const void * , void * , const size_t, const int, - const c_dpctl.DPCTLEventVectorRef) cpdef utils.dpnp_descriptor dpnp_absolute(utils.dpnp_descriptor x1): @@ -112,46 +106,6 @@ cpdef utils.dpnp_descriptor dpnp_absolute(utils.dpnp_descriptor x1): return result -cpdef utils.dpnp_descriptor dpnp_arctan2(utils.dpnp_descriptor x1_obj, - utils.dpnp_descriptor x2_obj, - object dtype=None, - utils.dpnp_descriptor out=None, - object where=True): - return call_fptr_2in_1out_strides(DPNP_FN_ARCTAN2_EXT, x1_obj, x2_obj, dtype, out, where, func_name="arctan2") - - -cpdef utils.dpnp_descriptor dpnp_around(utils.dpnp_descriptor x1, int decimals): - - cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype) - - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_AROUND_EXT, param1_type, param1_type) - - x1_obj = x1.get_array() - - # ceate result array with type given by FPTR data - cdef shape_type_c result_shape = x1.shape - cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, - kernel_data.return_type, - None, - device=x1_obj.sycl_device, - usm_type=x1_obj.usm_type, - sycl_queue=x1_obj.sycl_queue) - - result_sycl_queue = result.get_array().sycl_queue - - cdef c_dpctl.SyclQueue q = result_sycl_queue - cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() - - cdef ftpr_custom_around_1in_1out_t func = kernel_data.ptr - - cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, x1.get_data(), result.get_data(), x1.size, decimals, NULL) - - with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) - c_dpctl.DPCTLEvent_Delete(event_ref) - - return result - - cpdef utils.dpnp_descriptor dpnp_copysign(utils.dpnp_descriptor x1_obj, utils.dpnp_descriptor x2_obj, object dtype=None, @@ -298,7 +252,7 @@ cpdef utils.dpnp_descriptor dpnp_gradient(utils.dpnp_descriptor y1, int dx=1): # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils_py.create_output_descriptor_py(result_shape, - dpnp.float64, + dpnp.default_float_type(y1_obj.sycl_queue), None, device=y1_obj.sycl_device, usm_type=y1_obj.usm_type, @@ -388,7 +342,7 @@ cpdef tuple dpnp_modf(utils.dpnp_descriptor x1): cpdef utils.dpnp_descriptor dpnp_nancumprod(utils.dpnp_descriptor x1): - cur_x1 = dpnp_copy(x1).get_pyobj() + cur_x1 = x1.get_pyobj().copy() cur_x1_flatiter = cur_x1.flat @@ -401,7 +355,7 @@ cpdef utils.dpnp_descriptor dpnp_nancumprod(utils.dpnp_descriptor x1): cpdef utils.dpnp_descriptor dpnp_nancumsum(utils.dpnp_descriptor x1): - cur_x1 = dpnp_copy(x1).get_pyobj() + cur_x1 = x1.get_pyobj().copy() cur_x1_flatiter = cur_x1.flat @@ -453,14 +407,6 @@ cpdef utils.dpnp_descriptor dpnp_nansum(utils.dpnp_descriptor x1): return dpnp_sum(result) -cpdef utils.dpnp_descriptor dpnp_power(utils.dpnp_descriptor x1_obj, - utils.dpnp_descriptor x2_obj, - object dtype=None, - utils.dpnp_descriptor out=None, - object where=True): - return call_fptr_2in_1out_strides(DPNP_FN_POWER_EXT, x1_obj, x2_obj, dtype, out, where, func_name="power") - - cpdef utils.dpnp_descriptor dpnp_prod(utils.dpnp_descriptor x1, object axis=None, object dtype=None, diff --git a/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi b/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi index ec129a5285e6..55e7e9c8dc53 100644 --- a/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi @@ -36,14 +36,7 @@ and the rest of the library # NO IMPORTs here. All imports must be placed into main "dpnp_algo.pyx" file __all__ += [ - 'dpnp_arccos', - 'dpnp_arccosh', - 'dpnp_arcsin', - 'dpnp_arcsinh', - 'dpnp_arctan', - 'dpnp_arctanh', 'dpnp_cbrt', - 'dpnp_cosh', 'dpnp_degrees', 'dpnp_exp', 'dpnp_exp2', @@ -53,45 +46,14 @@ __all__ += [ 'dpnp_log2', 'dpnp_radians', 'dpnp_recip', - 'dpnp_sinh', - 'dpnp_tan', - 'dpnp_tanh', 'dpnp_unwrap' ] -cpdef utils.dpnp_descriptor dpnp_arccos(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_ARCCOS_EXT, x1) - - -cpdef utils.dpnp_descriptor dpnp_arccosh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_ARCCOSH_EXT, x1) - - -cpdef utils.dpnp_descriptor dpnp_arcsin(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_ARCSIN_EXT, x1, dtype=None, out=out, where=True, func_name='arcsin') - - -cpdef utils.dpnp_descriptor dpnp_arcsinh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_ARCSINH_EXT, x1) - - -cpdef utils.dpnp_descriptor dpnp_arctan(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_ARCTAN_EXT, x1, dtype=None, out=out, where=True, func_name='arctan') - - -cpdef utils.dpnp_descriptor dpnp_arctanh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_ARCTANH_EXT, x1) - - cpdef utils.dpnp_descriptor dpnp_cbrt(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_CBRT_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_cosh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_COSH_EXT, x1) - - cpdef utils.dpnp_descriptor dpnp_degrees(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_DEGREES_EXT, x1) @@ -128,18 +90,6 @@ cpdef utils.dpnp_descriptor dpnp_radians(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_RADIANS_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_sinh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_SINH_EXT, x1) - - -cpdef utils.dpnp_descriptor dpnp_tan(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_TAN_EXT, x1, dtype=None, out=out, where=True, func_name='tan') - - -cpdef utils.dpnp_descriptor dpnp_tanh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_TANH_EXT, x1) - - cpdef utils.dpnp_descriptor dpnp_unwrap(utils.dpnp_descriptor array1): result_type = dpnp.float64 diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index dd3a34369e3d..ffc1f20065c4 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -27,8 +27,8 @@ # ***************************************************************************** -import dpctl -import dpctl.tensor as dpt +from sys import platform + import dpctl.tensor._tensor_impl as ti from dpctl.tensor._elementwise_common import ( BinaryElementwiseFunc, @@ -42,13 +42,21 @@ __all__ = [ "check_nd_call_func", + "dpnp_acos", + "dpnp_acosh", "dpnp_add", + "dpnp_asin", + "dpnp_asinh", + "dpnp_atan", + "dpnp_atan2", + "dpnp_atanh", "dpnp_bitwise_and", "dpnp_bitwise_or", "dpnp_bitwise_xor", "dpnp_ceil", "dpnp_conj", "dpnp_cos", + "dpnp_cosh", "dpnp_divide", "dpnp_equal", "dpnp_floor", @@ -68,13 +76,22 @@ "dpnp_logical_or", "dpnp_logical_xor", "dpnp_multiply", + "dpnp_negative", "dpnp_not_equal", + "dpnp_power", + "dpnp_proj", "dpnp_remainder", "dpnp_right_shift", + "dpnp_round", + "dpnp_sign", + "dpnp_signbit", "dpnp_sin", + "dpnp_sinh", "dpnp_sqrt", "dpnp_square", "dpnp_subtract", + "dpnp_tan", + "dpnp_tanh", "dpnp_trunc", ] @@ -145,6 +162,114 @@ def check_nd_call_func( ) +_acos_docstring = """ +acos(x, out=None, order='K') + +Computes inverse cosine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise inverse cosine, in radians + and in the closed interval `[-pi/2, pi/2]`. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_acos(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_acos_to_call(sycl_queue, src, dst): + # call pybind11 extension for acos() function from OneMKL VM + return vmi._acos(sycl_queue, src, dst, depends) + return ti._acos(src, dst, sycl_queue, depends) + + +acos_func = UnaryElementwiseFunc( + "acos", ti._acos_result_type, _call_acos, _acos_docstring +) + + +def dpnp_acos(x, out=None, order="K"): + """ + Invokes acos() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for acos() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = acos_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + +_acosh_docstring = """ +acosh(x, out=None, order='K') + +Computes hyperbolic inverse cosine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise inverse hyperbolic cosine. + The data type of the returned array is determined by + the Type Promotion Rules. +""" + + +def _call_acosh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_acosh_to_call(sycl_queue, src, dst): + # call pybind11 extension for acosh() function from OneMKL VM + return vmi._acosh(sycl_queue, src, dst, depends) + return ti._acosh(src, dst, sycl_queue, depends) + + +acosh_func = UnaryElementwiseFunc( + "acosh", ti._acosh_result_type, _call_acosh, _acosh_docstring +) + + +def dpnp_acosh(x, out=None, order="K"): + """ + Invokes acosh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for acosh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = acosh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _add_docstring_ = """ add(x1, x2, out=None, order="K") @@ -204,6 +329,288 @@ def dpnp_add(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_asin_docstring = """ +asin(x, out=None, order='K') + +Computes inverse sine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise inverse sine, in radians + and in the closed interval `[-pi/2, pi/2]`. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_asin(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_asin_to_call(sycl_queue, src, dst): + # call pybind11 extension for asin() function from OneMKL VM + return vmi._asin(sycl_queue, src, dst, depends) + return ti._asin(src, dst, sycl_queue, depends) + + +asin_func = UnaryElementwiseFunc( + "asin", ti._asin_result_type, _call_asin, _asin_docstring +) + + +def dpnp_asin(x, out=None, order="K"): + """ + Invokes asin() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for asin() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = asin_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + +_asinh_docstring = """ +asinh(x, out=None, order='K') + +Computes inverse hyperbolic sine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise inverse hyperbolic sine. + The data type of the returned array is determined by + the Type Promotion Rules. +""" + + +def _call_asinh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_asinh_to_call(sycl_queue, src, dst): + # call pybind11 extension for asinh() function from OneMKL VM + return vmi._asinh(sycl_queue, src, dst, depends) + return ti._asinh(src, dst, sycl_queue, depends) + + +asinh_func = UnaryElementwiseFunc( + "asinh", ti._asinh_result_type, _call_asinh, _asinh_docstring +) + + +def dpnp_asinh(x, out=None, order="K"): + """ + Invokes asinh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for asinh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = asinh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + +_atan_docstring = """ +atan(x, out=None, order='K') + +Computes inverse tangent for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise inverse tangent, in radians + and in the closed interval `[-pi/2, pi/2]`. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_atan(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_atan_to_call(sycl_queue, src, dst): + # call pybind11 extension for atan() function from OneMKL VM + return vmi._atan(sycl_queue, src, dst, depends) + return ti._atan(src, dst, sycl_queue, depends) + + +atan_func = UnaryElementwiseFunc( + "atan", ti._atan_result_type, _call_atan, _atan_docstring +) + + +def dpnp_atan(x, out=None, order="K"): + """ + Invokes atan() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for atan() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = atan_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + +_atan2_docstring_ = """ +atan2(x1, x2, out=None, order="K") + +Calculates the inverse tangent of the quotient `x1_i/x2_i` for each element +`x1_i` of the input array `x1` with the respective element `x2_i` of the +input array `x2`. Each element-wise result is expressed in radians. + +Args: + x1 (dpnp.ndarray): + First input array, expected to have a real-valued floating-point + data type. + x2 (dpnp.ndarray): + Second input array, also expected to have a real-valued + floating-point data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. + Array have the correct shape and the expected data type. + order ("C","F","A","K", None, optional): + Memory layout of the newly output array, if parameter `out` is `None`. + Default: "K". +Returns: + dpnp.ndarray: + An array containing the inverse tangent of the quotient `x1`/`x2`. + The returned array must have a real-valued floating-point data type + determined by Type Promotion Rules. +""" + + +def _call_atan2(src1, src2, dst, sycl_queue, depends=None): + """A callback to register in BinaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_atan2_to_call(sycl_queue, src1, src2, dst): + # call pybind11 extension for atan2() function from OneMKL VM + return vmi._atan2(sycl_queue, src1, src2, dst, depends) + return ti._atan2(src1, src2, dst, sycl_queue, depends) + + +atan2_func = BinaryElementwiseFunc( + "atan2", + ti._atan2_result_type, + _call_atan2, + _atan2_docstring_, +) + + +def dpnp_atan2(x1, x2, out=None, order="K"): + """ + Invokes atan2() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for atan2() function. + """ + + # dpctl.tensor only works with usm_ndarray or scalar + x1_usm_or_scalar = dpnp.get_usm_ndarray_or_scalar(x1) + x2_usm_or_scalar = dpnp.get_usm_ndarray_or_scalar(x2) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = atan2_func( + x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order + ) + return dpnp_array._create_from_usm_ndarray(res_usm) + + +_atanh_docstring = """ +atanh(x, out=None, order='K') + +Computes hyperbolic inverse tangent for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise hyperbolic inverse tangent. + The data type of the returned array is determined by + the Type Promotion Rules. +""" + + +def _call_atanh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_atanh_to_call(sycl_queue, src, dst): + # call pybind11 extension for atanh() function from OneMKL VM + return vmi._atanh(sycl_queue, src, dst, depends) + return ti._atanh(src, dst, sycl_queue, depends) + + +atanh_func = UnaryElementwiseFunc( + "atanh", ti._atanh_result_type, _call_atanh, _atanh_docstring +) + + +def dpnp_atanh(x, out=None, order="K"): + """ + Invokes atanh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for atanh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = atanh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _bitwise_and_docstring_ = """ bitwise_and(x1, x2, out=None, order='K') @@ -403,7 +810,9 @@ def dpnp_ceil(x, out=None, order="K"): _cos_docstring = """ cos(x, out=None, order='K') + Computes cosine for each element `x_i` for input array `x`. + Args: x (dpnp.ndarray): Input array, expected to have numeric data type. @@ -420,6 +829,91 @@ def dpnp_ceil(x, out=None, order="K"): """ +def _call_cos(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_cos_to_call(sycl_queue, src, dst): + # call pybind11 extension for cos() function from OneMKL VM + return vmi._cos(sycl_queue, src, dst, depends) + return ti._cos(src, dst, sycl_queue, depends) + + +cos_func = UnaryElementwiseFunc( + "cos", ti._cos_result_type, _call_cos, _cos_docstring +) + + +def dpnp_cos(x, out=None, order="K"): + """ + Invokes cos() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for cos() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = cos_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + +_cosh_docstring = """ +cosh(x, out=None, order='K') + +Computes hyperbolic cosine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise hyperbolic cosine. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_cosh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_cosh_to_call(sycl_queue, src, dst): + # call pybind11 extension for cosh() function from OneMKL VM + return vmi._cosh(sycl_queue, src, dst, depends) + return ti._cosh(src, dst, sycl_queue, depends) + + +cosh_func = UnaryElementwiseFunc( + "cosh", ti._cosh_result_type, _call_cosh, _cosh_docstring +) + + +def dpnp_cosh(x, out=None, order="K"): + """ + Invokes cosh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for cosh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = cosh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _conj_docstring = """ conj(x, out=None, order='K') @@ -458,36 +952,6 @@ def _call_conj(src, dst, sycl_queue, depends=None): ) -def dpnp_cos(x, out=None, order="K"): - """ - Invokes cos() function from pybind11 extension of OneMKL VM if possible. - - Otherwise fully relies on dpctl.tensor implementation for cos() function. - - """ - - def _call_cos(src, dst, sycl_queue, depends=None): - """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" - - if depends is None: - depends = [] - - if vmi._mkl_cos_to_call(sycl_queue, src, dst): - # call pybind11 extension for cos() function from OneMKL VM - return vmi._cos(sycl_queue, src, dst, depends) - return ti._cos(src, dst, sycl_queue, depends) - - # dpctl.tensor only works with usm_ndarray - x1_usm = dpnp.get_usm_ndarray(x) - out_usm = None if out is None else dpnp.get_usm_ndarray(out) - - func = UnaryElementwiseFunc( - "cos", ti._cos_result_type, _call_cos, _cos_docstring - ) - res_usm = func(x1_usm, out=out_usm, order=order) - return dpnp_array._create_from_usm_ndarray(res_usm) - - def dpnp_conj(x, out=None, order="K"): """ Invokes conj() function from pybind11 extension of OneMKL VM if possible. @@ -530,32 +994,12 @@ def _call_divide(src1, src2, dst, sycl_queue, depends=None): """A callback to register in BinaryElementwiseFunc class of dpctl.tensor""" if depends is None: - depends = [] - - if vmi._mkl_div_to_call(sycl_queue, src1, src2, dst): - # call pybind11 extension for div() function from OneMKL VM - return vmi._div(sycl_queue, src1, src2, dst, depends) - return ti._divide(src1, src2, dst, sycl_queue, depends) - - -def _call_divide_inplace(lhs, rhs, sycl_queue, depends=None): - """In place workaround until dpctl.tensor provides the functionality.""" - - if depends is None: - depends = [] - - # allocate temporary memory for out array - out = dpt.empty_like(lhs, dtype=dpnp.result_type(lhs.dtype, rhs.dtype)) - - # call a general callback - div_ht_, div_ev_ = _call_divide(lhs, rhs, out, sycl_queue, depends) + depends = [] - # store the result into left input array and return events - cp_ht_, cp_ev_ = ti._copy_usm_ndarray_into_usm_ndarray( - src=out, dst=lhs, sycl_queue=sycl_queue, depends=[div_ev_] - ) - dpctl.SyclEvent.wait_for([div_ht_]) - return (cp_ht_, cp_ev_) + if vmi._mkl_div_to_call(sycl_queue, src1, src2, dst): + # call pybind11 extension for div() function from OneMKL VM + return vmi._div(sycl_queue, src1, src2, dst, depends) + return ti._divide(src1, src2, dst, sycl_queue, depends) divide_func = BinaryElementwiseFunc( @@ -563,7 +1007,6 @@ def _call_divide_inplace(lhs, rhs, sycl_queue, depends=None): ti._divide_result_type, _call_divide, _divide_docstring_, - _call_divide_inplace, ) @@ -995,8 +1438,8 @@ def dpnp_isnan(x, out=None, order="K"): """ -leftt_shift_func = BinaryElementwiseFunc( - "bitwise_leftt_shift", +left_shift_func = BinaryElementwiseFunc( + "bitwise_left_shift", ti._bitwise_left_shift_result_type, ti._bitwise_left_shift, _left_shift_docstring_, @@ -1011,7 +1454,7 @@ def dpnp_left_shift(x1, x2, out=None, order="K"): x2_usm_or_scalar = dpnp.get_usm_ndarray_or_scalar(x2) out_usm = None if out is None else dpnp.get_usm_ndarray(out) - res_usm = leftt_shift_func( + res_usm = left_shift_func( x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order ) return dpnp_array._create_from_usm_ndarray(res_usm) @@ -1478,6 +1921,106 @@ def dpnp_not_equal(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_power_docstring_ = """ +power(x1, x2, out=None, order="K") + +Calculates `x1_i` raised to `x2_i` for each element `x1_i` of the input array +`x1` with the respective element `x2_i` of the input array `x2`. + +Args: + x1 (dpnp.ndarray): + First input array, expected to have numeric data type. + x2 (dpnp.ndarray): + Second input array, also expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", None, optional): + Output array, if parameter `out` is `None`. + Default: "K". +Returns: + dpnp.ndarray: + An array containing the result of element-wise of raising each element + to a specified power. + The data type of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_pow(src1, src2, dst, sycl_queue, depends=None): + """A callback to register in BinaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + # TODO: remove this check when OneMKL is fixed on Windows + is_win = platform.startswith("win") + + if not is_win and vmi._mkl_pow_to_call(sycl_queue, src1, src2, dst): + # call pybind11 extension for pow() function from OneMKL VM + return vmi._pow(sycl_queue, src1, src2, dst, depends) + return ti._pow(src1, src2, dst, sycl_queue, depends) + + +pow_func = BinaryElementwiseFunc( + "pow", ti._pow_result_type, _call_pow, _power_docstring_ +) + + +def dpnp_power(x1, x2, out=None, order="K"): + """ + Invokes pow() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for pow() function. + """ + + # dpctl.tensor only works with usm_ndarray or scalar + x1_usm_or_scalar = dpnp.get_usm_ndarray_or_scalar(x1) + x2_usm_or_scalar = dpnp.get_usm_ndarray_or_scalar(x2) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = pow_func( + x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order + ) + return dpnp_array._create_from_usm_ndarray(res_usm) + + +_proj_docstring = """ +proj(x, out=None, order="K") + +Computes projection of each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. + Array have the correct shape and the expected data type. + order ("C","F","A","K", optional): + Memory layout of the newly output array, if parameter `out` is `None`. + Default: "K". +Returns: + dpnp.ndarray: + An array containing the element-wise projection. + The returned array has the same data type as `x`. +""" + + +proj_func = UnaryElementwiseFunc( + "proj", ti._proj_result_type, ti._proj, _proj_docstring +) + + +def dpnp_proj(x, out=None, order="K"): + """Invokes proj() from dpctl.tensor implementation for proj() function.""" + + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = proj_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _remainder_docstring_ = """ remainder(x1, x2, out=None, order='K') Calculates the remainder of division for each element `x1_i` of the input array @@ -1568,6 +2111,57 @@ def dpnp_right_shift(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_round_docstring = """ +round(x, out=None, order='K') +Rounds each element `x_i` of the input array `x` to +the nearest integer-valued number. +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise rounded value. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_round(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_round_to_call(sycl_queue, src, dst): + # call pybind11 extension for round() function from OneMKL VM + return vmi._round(sycl_queue, src, dst, depends) + return ti._round(src, dst, sycl_queue, depends) + + +round_func = UnaryElementwiseFunc( + "round", ti._round_result_type, _call_round, _round_docstring +) + + +def dpnp_round(x, out=None, order="K"): + """ + Invokes round() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for round() function. + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = round_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _sign_docstring = """ sign(x, out=None, order="K") @@ -1613,9 +2207,49 @@ def dpnp_sign(x, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_signbit_docstring = """ +signbit(x, out=None, order="K") + +Computes an indication of whether the sign bit of each element `x_i` of +input array `x` is set. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. + Array have the correct shape and the expected data type. + order ("C","F","A","K", optional): + Memory layout of the newly output array, if parameter `out` is `None`. + Default: "K". +Returns: + dpnp.ndarray: + An array containing the element-wise results. The returned array + must have a data type of `bool`. +""" + + +signbit_func = UnaryElementwiseFunc( + "signbit", ti._signbit_result_type, ti._signbit, _signbit_docstring +) + + +def dpnp_signbit(x, out=None, order="K"): + """Invokes signbit() from dpctl.tensor implementation for signbit() function.""" + + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = signbit_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _sin_docstring = """ sin(x, out=None, order='K') + Computes sine for each element `x_i` of input array `x`. + Args: x (dpnp.ndarray): Input array, expected to have numeric data type. @@ -1664,6 +2298,59 @@ def dpnp_sin(x, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_sinh_docstring = """ +sinh(x, out=None, order='K') + +Computes hyperbolic sine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise hyperbolic sine. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_sinh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_sinh_to_call(sycl_queue, src, dst): + # call pybind11 extension for sinh() function from OneMKL VM + return vmi._sinh(sycl_queue, src, dst, depends) + return ti._sinh(src, dst, sycl_queue, depends) + + +sinh_func = UnaryElementwiseFunc( + "sinh", ti._sinh_result_type, _call_sinh, _sinh_docstring +) + + +def dpnp_sinh(x, out=None, order="K"): + """ + Invokes sinh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for sinh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = sinh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _sqrt_docstring_ = """ sqrt(x, out=None, order='K') Computes the non-negative square-root for each element `x_i` for input array `x`. @@ -1844,6 +2531,112 @@ def dpnp_subtract(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_tan_docstring = """ +tan(x, out=None, order='K') + +Computes tangent for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise tangent. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_tan(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_tan_to_call(sycl_queue, src, dst): + # call pybind11 extension for tan() function from OneMKL VM + return vmi._tan(sycl_queue, src, dst, depends) + return ti._tan(src, dst, sycl_queue, depends) + + +tan_func = UnaryElementwiseFunc( + "tan", ti._tan_result_type, _call_tan, _tan_docstring +) + + +def dpnp_tan(x, out=None, order="K"): + """ + Invokes tan() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for tan() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = tan_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + +_tanh_docstring = """ +tanh(x, out=None, order='K') + +Computes hyperbolic tangent for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise hyperbolic tangent. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_tanh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_tanh_to_call(sycl_queue, src, dst): + # call pybind11 extension for tanh() function from OneMKL VM + return vmi._tanh(sycl_queue, src, dst, depends) + return ti._tanh(src, dst, sycl_queue, depends) + + +tanh_func = UnaryElementwiseFunc( + "tanh", ti._tanh_result_type, _call_tanh, _tanh_docstring +) + + +def dpnp_tanh(x, out=None, order="K"): + """ + Invokes tanh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for tanh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = tanh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _trunc_docstring = """ trunc(x, out=None, order='K') diff --git a/dpnp/dpnp_array.py b/dpnp/dpnp_array.py index 717cb9c4f526..24bc3a03a5e8 100644 --- a/dpnp/dpnp_array.py +++ b/dpnp/dpnp_array.py @@ -168,7 +168,15 @@ def __complex__(self): return self._array_obj.__complex__() # '__contains__', - # '__copy__', + + def __copy__(self): + """ + Used if :func:`copy.copy` is called on an array. Returns a copy of the array. + + Equivalent to ``a.copy(order="K")``. + """ + return self.copy(order="K") + # '__deepcopy__', # '__delattr__', # '__delitem__', @@ -188,7 +196,10 @@ def __eq__(self, other): def __float__(self): return self._array_obj.__float__() - # '__floordiv__', + def __floordiv__(self, other): + """Return self//value.""" + return dpnp.floor_divide(self, other) + # '__format__', def __ge__(self, other): @@ -227,7 +238,10 @@ def __iand__(self, other): dpnp.bitwise_and(self, other, out=self) return self - # '__ifloordiv__', + def __ifloordiv__(self, other): + """Return self//=value.""" + dpnp.floor_divide(self, other, out=self) + return self def __ilshift__(self, other): """Return self<<=value.""" @@ -235,7 +249,11 @@ def __ilshift__(self, other): return self # '__imatmul__', - # '__imod__', + + def __imod__(self, other): + """Return self%=value.""" + dpnp.remainder(self, other, out=self) + return self def __imul__(self, other): """Return self*=value.""" @@ -345,7 +363,8 @@ def __rand__(self, other): def __repr__(self): return dpt.usm_ndarray_repr(self._array_obj, prefix="array") - # '__rfloordiv__', + def __rfloordiv__(self, other): + return dpnp.floor_divide(self, other) def __rlshift__(self, other): return dpnp.left_shift(other, self) @@ -640,7 +659,47 @@ def conjugate(self): else: return dpnp.conjugate(self) - # 'copy', + def copy(self, order="C"): + """ + Return a copy of the array. + + Returns + ------- + out : dpnp.ndarray + A copy of the array. + + See also + -------- + :obj:`dpnp.copy` : Similar function with different default behavior + :obj:`dpnp.copyto` : Copies values from one array to another. + + Notes + ----- + This function is the preferred method for creating an array copy. The + function :func:`dpnp.copy` is similar, but it defaults to using order 'K'. + + Examples + -------- + >>> import dpnp as np + >>> x = np.array([[1, 2, 3], [4, 5, 6]], order='F') + >>> y = x.copy() + >>> x.fill(0) + + >>> x + array([[0, 0, 0], + [0, 0, 0]]) + + >>> y + array([[1, 2, 3], + [4, 5, 6]]) + + >>> y.flags['C_CONTIGUOUS'] + True + + """ + + return dpnp.copy(self, order=order) + # 'ctypes', # 'cumprod', @@ -904,8 +963,17 @@ def prod( return dpnp.prod(self, axis, dtype, out, keepdims, initial, where) - # 'ptp', - # 'put', + # 'ptp' + + def put(self, indices, vals, /, *, axis=None, mode="wrap"): + """ + Puts values of an array into another array along a given axis. + + For full documentation refer to :obj:`numpy.put`. + """ + + return dpnp.put(self, indices, vals, axis=axis, mode=mode) + # 'ravel', # 'real', # 'repeat', @@ -1059,7 +1127,14 @@ def sum( where=where, ) - # 'swapaxes', + def swapaxes(self, axis1, axis2): + """ + Interchange two axes of an array. + + For full documentation refer to :obj:`numpy.swapaxes`. + """ + + return dpnp.swapaxes(self, axis1=axis1, axis2=axis2) def take(self, indices, /, *, axis=None, out=None, mode="wrap"): """ diff --git a/dpnp/dpnp_container.py b/dpnp/dpnp_container.py index 08522ceaafa6..faf14c3e97bb 100644 --- a/dpnp/dpnp_container.py +++ b/dpnp/dpnp_container.py @@ -43,6 +43,7 @@ __all__ = [ "arange", "asarray", + "copy", "empty", "eye", "full", @@ -135,6 +136,15 @@ def asarray( return dpnp_array(array_obj.shape, buffer=array_obj, order=order) +def copy(x1, /, *, order="K"): + """Creates `dpnp_array` as a copy of given instance of input array.""" + if order is None: + order = "K" + + array_obj = dpt.copy(dpnp.get_usm_ndarray(x1), order=order) + return dpnp_array(array_obj.shape, buffer=array_obj, order="K") + + def empty( shape, *, @@ -264,9 +274,7 @@ def meshgrid(*xi, indexing="xy"): """Creates list of `dpnp_array` coordinate matrices from vectors.""" if len(xi) == 0: return [] - arrays = tuple( - x.get_array() if isinstance(x, dpnp_array) else x for x in xi - ) + arrays = tuple(dpnp.get_usm_ndarray(x) for x in xi) arrays_obj = dpt.meshgrid(*arrays, indexing=indexing) return [ dpnp_array._create_from_usm_ndarray(array_obj) @@ -304,17 +312,13 @@ def ones( def tril(x1, /, *, k=0): """Creates `dpnp_array` as lower triangular part of an input array.""" - array_obj = dpt.tril( - x1.get_array() if isinstance(x1, dpnp_array) else x1, k - ) + array_obj = dpt.tril(dpnp.get_usm_ndarray(x1), k) return dpnp_array(array_obj.shape, buffer=array_obj, order="K") def triu(x1, /, *, k=0): """Creates `dpnp_array` as upper triangular part of an input array.""" - array_obj = dpt.triu( - x1.get_array() if isinstance(x1, dpnp_array) else x1, k - ) + array_obj = dpt.triu(dpnp.get_usm_ndarray(x1), k) return dpnp_array(array_obj.shape, buffer=array_obj, order="K") diff --git a/dpnp/dpnp_iface.py b/dpnp/dpnp_iface.py index 066b28e8e2cd..c131bf733ce1 100644 --- a/dpnp/dpnp_iface.py +++ b/dpnp/dpnp_iface.py @@ -66,6 +66,7 @@ "get_normalized_queue_device", "get_usm_ndarray", "get_usm_ndarray_or_scalar", + "is_supported_array_or_scalar", "is_supported_array_type", ] @@ -453,6 +454,28 @@ def get_usm_ndarray_or_scalar(a): return a if isscalar(a) else get_usm_ndarray(a) +def is_supported_array_or_scalar(a): + """ + Return ``True`` if `a` is a scalar or an array of either + :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray` type, + ``False`` otherwise. + + Parameters + ---------- + a : {scalar, dpnp_array, usm_ndarray} + An input scalar or an array to check the type of. + + Returns + ------- + out : bool + ``True`` if input `a` is a scalar or an array of supported type, + ``False`` otherwise. + + """ + + return isscalar(a) or is_supported_array_type(a) + + def is_supported_array_type(a): """ Return ``True`` if an array of either type :class:`dpnp.ndarray` @@ -460,7 +483,7 @@ def is_supported_array_type(a): Parameters ---------- - a : array + a : {dpnp_array, usm_ndarray} An input array to check the type. Returns diff --git a/dpnp/dpnp_iface_arraycreation.py b/dpnp/dpnp_iface_arraycreation.py index 549920066ac0..22e7c1e0d9d1 100644 --- a/dpnp/dpnp_iface_arraycreation.py +++ b/dpnp/dpnp_iface_arraycreation.py @@ -110,7 +110,7 @@ def arange( Returns ------- - arange : :obj:`dpnp.ndarray` + out : dpnp.ndarray The 1-D array containing evenly spaced values. Limitations @@ -151,7 +151,7 @@ def arange( def array( - x, + a, dtype=None, *, copy=True, @@ -168,6 +168,11 @@ def array( For full documentation refer to :obj:`numpy.array`. + Returns + ------- + out : dpnp.ndarray + An array object satisfying the specified requirements. + Limitations ----------- Parameter `subok` is supported only with default value ``False``. @@ -228,7 +233,7 @@ def array( copy = None return dpnp_container.asarray( - x, + a, dtype=dtype, copy=copy, order=order, @@ -239,7 +244,7 @@ def array( def asanyarray( - x, + a, dtype=None, order=None, *, @@ -249,10 +254,15 @@ def asanyarray( sycl_queue=None, ): """ - Convert the input to an ndarray, but pass ndarray subclasses through. + Convert the input to an :class:`dpnp.ndarray`. For full documentation refer to :obj:`numpy.asanyarray`. + Returns + ------- + out : dpnp.ndarray + Array interpretation of `a`. + Limitations ----------- Parameter `like` is supported only with default value ``None``. @@ -286,7 +296,7 @@ def asanyarray( ) return asarray( - x, + a, dtype=dtype, order=order, device=device, @@ -296,7 +306,7 @@ def asanyarray( def asarray( - x, + a, dtype=None, order=None, like=None, @@ -309,6 +319,12 @@ def asarray( For full documentation refer to :obj:`numpy.asarray`. + Returns + ------- + out : dpnp.ndarray + Array interpretation of `a`. No copy is performed if the input + is already an ndarray with matching dtype and order. + Limitations ----------- Parameter `like` is supported only with default value ``None``. @@ -342,7 +358,7 @@ def asarray( ) return dpnp_container.asarray( - x, + a, dtype=dtype, order=order, device=device, @@ -352,13 +368,19 @@ def asarray( def ascontiguousarray( - x, dtype=None, *, like=None, device=None, usm_type=None, sycl_queue=None + a, dtype=None, *, like=None, device=None, usm_type=None, sycl_queue=None ): """ Return a contiguous array (ndim >= 1) in memory (C order). For full documentation refer to :obj:`numpy.ascontiguousarray`. + Returns + ------- + out : dpnp.ndarray + Contiguous array of same shape and content as `a`, with type `dtype` + if specified. + Limitations ----------- Parameter `like` is supported only with default value ``None``. @@ -407,11 +429,11 @@ def ascontiguousarray( ) # at least 1-d array has to be returned - if x.ndim == 0: - x = [x] + if a.ndim == 0: + a = [a] return asarray( - x, + a, dtype=dtype, order="C", device=device, @@ -421,13 +443,18 @@ def ascontiguousarray( def asfortranarray( - x, dtype=None, *, like=None, device=None, usm_type=None, sycl_queue=None + a, dtype=None, *, like=None, device=None, usm_type=None, sycl_queue=None ): """ Return an array (ndim >= 1) laid out in Fortran order in memory. For full documentation refer to :obj:`numpy.asfortranarray`. + Returns + ------- + out : dpnp.ndarray + The input `a` in Fortran, or column-major, order. + Limitations ----------- Parameter `like` is supported only with default value ``None``. @@ -479,11 +506,11 @@ def asfortranarray( ) # at least 1-d array has to be returned - if x.ndim == 0: - x = [x] + if a.ndim == 0: + a = [a] return asarray( - x, + a, dtype=dtype, order="F", device=device, @@ -492,7 +519,7 @@ def asfortranarray( ) -def copy(x1, order="K", subok=False): +def copy(a, order="K", subok=False): """ Return an array copy of the given object. @@ -500,35 +527,53 @@ def copy(x1, order="K", subok=False): Limitations ----------- - Parameter ``order`` is supported only with default value ``"C"``. - Parameter ``subok`` is supported only with default value ``False``. + Parameter `subok` is supported only with default value ``False``. + Otherwise, the function raises `ValueError` exception. + + Returns + ------- + out : dpnp.ndarray + Array interpretation of `a`. + + See Also + -------- + :obj:`dpnp.ndarray.copy` : Preferred method for creating an array copy + + Notes + ----- + This is equivalent to: + + >>> dpnp.array(a, copy=True) Examples -------- + Create an array `x`, with a reference `y` and a copy `z`: + >>> import dpnp as np >>> x = np.array([1, 2, 3]) >>> y = x >>> z = np.copy(x) + + Note that, when we modify `x`, `y` will change, but not `z`: + >>> x[0] = 10 >>> x[0] == y[0] - True + array(True) >>> x[0] == z[0] - False + array(False) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False - ) - if x1_desc: - if order != "K": - pass - elif subok: - pass - else: - return dpnp_copy(x1_desc).get_pyobj() + if subok is not False: + raise ValueError( + "Keyword argument `subok` is supported only with " + f"default value ``False``, but got {subok}" + ) + + if dpnp.is_supported_array_type(a): + return dpnp_container.copy(a, order=order) - return call_origin(numpy.copy, x1, order, subok) + return array(a, order=order, subok=subok, copy=True) def diag(x1, k=0): diff --git a/dpnp/dpnp_iface_bitwise.py b/dpnp/dpnp_iface_bitwise.py index eded1984809b..bac53b2cbc6a 100644 --- a/dpnp/dpnp_iface_bitwise.py +++ b/dpnp/dpnp_iface_bitwise.py @@ -90,7 +90,7 @@ def bitwise_and( Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Data type of input arrays `x1` and `x2` has to be an integer or boolean data type. @@ -160,7 +160,7 @@ def bitwise_or( Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Data type of input arrays `x1` and `x2` has to be an integer or boolean data type. @@ -225,7 +225,7 @@ def bitwise_xor( Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Data type of input arrays `x1` and `x2` has to be an integer or boolean data type. @@ -286,7 +286,7 @@ def invert( Parameter `x` is supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Data type of input array `x` has to be an integer data type. @@ -358,7 +358,7 @@ def left_shift( Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input data is supported as integer only. @@ -421,7 +421,7 @@ def right_shift( Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input data is supported as integer only. diff --git a/dpnp/dpnp_iface_indexing.py b/dpnp/dpnp_iface_indexing.py index e378589b863b..0c19127831a1 100644 --- a/dpnp/dpnp_iface_indexing.py +++ b/dpnp/dpnp_iface_indexing.py @@ -417,34 +417,81 @@ def place(x, mask, vals, /): return call_origin(numpy.place, x, mask, vals, dpnp_inplace=True) -def put(x1, ind, v, mode="raise"): +def put(a, indices, vals, /, *, axis=None, mode="wrap"): """ - Replaces specified elements of an array with given values. + Puts values of an array into another array along a given axis. For full documentation refer to :obj:`numpy.put`. Limitations ----------- - Input array is supported as :obj:`dpnp.ndarray`. - Not supported parameter mode. + Parameters `a` and `indices` are supported either as :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. + Parameter `indices` is supported as 1-D array of integer data type. + Parameter `vals` must be broadcastable to the shape of `indices` + and has the same data type as `a` if it is as :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. + Parameter `mode` is supported with ``wrap``, the default, and ``clip`` values. + Parameter `axis` is supported as integer only. + Otherwise the function will be executed sequentially on CPU. + + See Also + -------- + :obj:`dpnp.putmask` : Changes elements of an array based on conditional and input values. + :obj:`dpnp.place` : Change elements of an array based on conditional and input values. + :obj:`dpnp.put_along_axis` : Put values into the destination array by matching 1d index and data slices. + + Notes + ----- + In contrast to :obj:`numpy.put` `wrap` mode which wraps indices around the array for cyclic operations, + :obj:`dpnp.put` `wrap` mode clamps indices to a fixed range within the array boundaries (-n <= i < n). + + Examples + -------- + >>> import dpnp as np + >>> x = np.arange(5) + >>> indices = np.array([0, 1]) + >>> np.put(x, indices, [-44, -55]) + >>> x + array([-44, -55, 2, 3, 4]) + + >>> x = np.arange(5) + >>> indices = np.array([22]) + >>> np.put(x, indices, -5, mode='clip') + >>> x + array([ 0, 1, 2, 3, -5]) + """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False - ) - if x1_desc: - if mode != "raise": + if dpnp.is_supported_array_type(a) and dpnp.is_supported_array_type( + indices + ): + if indices.ndim != 1 or not dpnp.issubdtype( + indices.dtype, dpnp.integer + ): pass - elif type(ind) is not type(v): + elif mode not in ("clip", "wrap"): pass - elif ( - numpy.max(ind) >= x1_desc.size or numpy.min(ind) + x1_desc.size < 0 - ): + elif axis is not None and not isinstance(axis, int): + raise TypeError(f"`axis` must be of integer type, got {type(axis)}") + # TODO: remove when #1382(dpctl) is solved + elif dpnp.is_supported_array_type(vals) and a.dtype != vals.dtype: pass else: - return dpnp_put(x1_desc, ind, v) + if axis is None and a.ndim > 1: + a = dpnp.reshape(a, -1) + dpt_array = dpnp.get_usm_ndarray(a) + dpt_indices = dpnp.get_usm_ndarray(indices) + dpt_vals = ( + dpnp.get_usm_ndarray(vals) + if isinstance(vals, dpnp_array) + else vals + ) + return dpt.put( + dpt_array, dpt_indices, dpt_vals, axis=axis, mode=mode + ) - return call_origin(numpy.put, x1, ind, v, mode, dpnp_inplace=True) + return call_origin(numpy.put, a, indices, vals, mode, dpnp_inplace=True) def put_along_axis(x1, indices, values, axis): @@ -557,7 +604,7 @@ def take(x, indices, /, *, axis=None, out=None, mode="wrap"): or :class:`dpctl.tensor.usm_ndarray`. Parameter `indices` is supported as 1-D array of integer data type. Parameter `out` is supported only with default value. - Parameter `mode` is supported with ``wrap``(default) and ``clip`` mode. + Parameter `mode` is supported with ``wrap``, the default, and ``clip`` values. Providing parameter `axis` is optional when `x` is a 1-D array. Otherwise the function will be executed sequentially on CPU. diff --git a/dpnp/dpnp_iface_linearalgebra.py b/dpnp/dpnp_iface_linearalgebra.py index 70e2a8048c0e..3e3e92d82450 100644 --- a/dpnp/dpnp_iface_linearalgebra.py +++ b/dpnp/dpnp_iface_linearalgebra.py @@ -358,14 +358,15 @@ def outer(x1, x2, out=None): [1, 2, 3]]) """ + x1_is_scalar = dpnp.isscalar(x1) x2_is_scalar = dpnp.isscalar(x2) if x1_is_scalar and x2_is_scalar: pass - elif not (x1_is_scalar or dpnp.is_supported_array_type(x1)): + elif not dpnp.is_supported_array_or_scalar(x1): pass - elif not (x2_is_scalar or dpnp.is_supported_array_type(x2)): + elif not dpnp.is_supported_array_or_scalar(x2): pass else: x1_in = ( diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index d867bbbd6609..c4704d09f9a0 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -94,9 +94,9 @@ def all(x, /, axis=None, out=None, keepdims=False, *, where=True): Returns ------- - dpnp.ndarray - An array with a data type of `bool` - containing the results of the logical AND reduction. + out : dpnp.ndarray + An array with a data type of `bool` + containing the results of the logical AND reduction. Limitations ----------- @@ -152,42 +152,94 @@ def all(x, /, axis=None, out=None, keepdims=False, *, where=True): ) -def allclose(x1, x2, rtol=1.0e-5, atol=1.0e-8, **kwargs): +def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, **kwargs): """ Returns True if two arrays are element-wise equal within a tolerance. For full documentation refer to :obj:`numpy.allclose`. + Returns + ------- + out : dpnp.ndarray + A boolean 0-dim array. If its value is ``True``, + two arrays are element-wise equal within a tolerance. + Limitations ----------- - Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar. + Parameters `a` and `b` are supported either as :class:`dpnp.ndarray`, + :class:`dpctl.tensor.usm_ndarray` or scalars, but both `a` and `b` + can not be scalars at the same time. Keyword argument `kwargs` is currently unsupported. Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameters `rtol` and `atol` are supported as scalars. Otherwise + ``TypeError`` exeption will be raised. + Input array data types are limited by supported integer and + floating DPNP :ref:`Data types`. + + See Also + -------- + :obj:`dpnp.isclose` : Test whether two arrays are element-wise equal. + :obj:`dpnp.all` : Test whether all elements evaluate to True. + :obj:`dpnp.any` : Test whether any element evaluates to True. + :obj:`dpnp.equal` : Return (x1 == x2) element-wise. Examples -------- >>> import dpnp as np - >>> np.allclose([1e10,1e-7], [1.00001e10,1e-8]) - >>> False + >>> a = np.array([1e10, 1e-7]) + >>> b = np.array([1.00001e10, 1e-8]) + >>> np.allclose(a, b) + array([False]) + + >>> a = np.array([1.0, np.nan]) + >>> b = np.array([1.0, np.nan]) + >>> np.allclose(a, b) + array([False]) + + >>> a = np.array([1.0, np.inf]) + >>> b = np.array([1.0, np.inf]) + >>> np.allclose(a, b) + array([ True]) """ - rtol_is_scalar = dpnp.isscalar(rtol) - atol_is_scalar = dpnp.isscalar(atol) - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - x2_desc = dpnp.get_dpnp_descriptor(x2, copy_when_nondefault_queue=False) + if dpnp.isscalar(a) and dpnp.isscalar(b): + # at least one of inputs has to be an array + pass + elif not ( + dpnp.is_supported_array_or_scalar(a) + and dpnp.is_supported_array_or_scalar(b) + ): + pass + elif kwargs: + pass + else: + if not dpnp.isscalar(rtol): + raise TypeError( + "An argument `rtol` must be a scalar, but got {}".format( + type(rtol) + ) + ) + elif not dpnp.isscalar(atol): + raise TypeError( + "An argument `atol` must be a scalar, but got {}".format( + type(atol) + ) + ) - if x1_desc and x2_desc and not kwargs: - if not rtol_is_scalar or not atol_is_scalar: - pass - else: - result_obj = dpnp_allclose(x1_desc, x2_desc, rtol, atol).get_pyobj() - result = dpnp.convert_single_elem_array_to_scalar(result_obj) + if dpnp.isscalar(a): + a = dpnp.full_like(b, fill_value=a) + elif dpnp.isscalar(b): + b = dpnp.full_like(a, fill_value=b) + elif a.shape != b.shape: + a, b = dpt.broadcast_arrays(a.get_array(), b.get_array()) - return result + a_desc = dpnp.get_dpnp_descriptor(a, copy_when_nondefault_queue=False) + b_desc = dpnp.get_dpnp_descriptor(b, copy_when_nondefault_queue=False) + if a_desc and b_desc: + return dpnp_allclose(a_desc, b_desc, rtol, atol).get_pyobj() - return call_origin(numpy.allclose, x1, x2, rtol=rtol, atol=atol, **kwargs) + return call_origin(numpy.allclose, a, b, rtol=rtol, atol=atol, **kwargs) def any(x, /, axis=None, out=None, keepdims=False, *, where=True): @@ -198,7 +250,7 @@ def any(x, /, axis=None, out=None, keepdims=False, *, where=True): Returns ------- - dpnp.ndarray + out : dpnp.ndarray An array with a data type of `bool` containing the results of the logical OR reduction. @@ -276,7 +328,7 @@ def equal( Returns ------- out : dpnp.ndarray - Output array of bool type, element-wise comparison of `x1` and `x2`. + Output array of bool type, element-wise comparison of `x1` and `x2`. Limitations ----------- @@ -352,7 +404,7 @@ def greater( Returns ------- out : dpnp.ndarray - Output array of bool type, element-wise comparison of `x1` and `x2`. + Output array of bool type, element-wise comparison of `x1` and `x2`. Limitations ----------- @@ -422,7 +474,7 @@ def greater_equal( Returns ------- out : dpnp.ndarray - Output array of bool type, element-wise comparison of `x1` and `x2`. + Output array of bool type, element-wise comparison of `x1` and `x2`. Limitations ----------- @@ -678,7 +730,7 @@ def less( Returns ------- out : dpnp.ndarray - Output array of bool type, element-wise comparison of `x1` and `x2`. + Output array of bool type, element-wise comparison of `x1` and `x2`. Limitations ----------- @@ -748,7 +800,7 @@ def less_equal( Returns ------- out : dpnp.ndarray - Output array of bool type, element-wise comparison of `x1` and `x2`. + Output array of bool type, element-wise comparison of `x1` and `x2`. Limitations ----------- @@ -818,8 +870,8 @@ def logical_and( Returns ------- out : dpnp.ndarray - Boolean result of the logical AND operation applied to the elements - of `x1` and `x2`; the shape is determined by broadcasting. + Boolean result of the logical AND operation applied to the elements + of `x1` and `x2`; the shape is determined by broadcasting. Limitations ----------- @@ -891,8 +943,8 @@ def logical_not( Returns ------- out : dpnp.ndarray - Boolean result with the same shape as `x` of the NOT operation - on elements of `x`. + Boolean result with the same shape as `x` of the NOT operation + on elements of `x`. Limitations ----------- @@ -953,8 +1005,8 @@ def logical_or( Returns ------- out : dpnp.ndarray - Boolean result of the logical OR operation applied to the elements - of `x1` and `x2`; the shape is determined by broadcasting. + Boolean result of the logical OR operation applied to the elements + of `x1` and `x2`; the shape is determined by broadcasting. Limitations ----------- @@ -1027,8 +1079,8 @@ def logical_xor( Returns ------- out : dpnp.ndarray - Boolean result of the logical XOR operation applied to the elements - of `x1` and `x2`; the shape is determined by broadcasting. + Boolean result of the logical XOR operation applied to the elements + of `x1` and `x2`; the shape is determined by broadcasting. Limitations ----------- diff --git a/dpnp/dpnp_iface_manipulation.py b/dpnp/dpnp_iface_manipulation.py index e29e6b0203ea..f71aa751cb2c 100644 --- a/dpnp/dpnp_iface_manipulation.py +++ b/dpnp/dpnp_iface_manipulation.py @@ -42,11 +42,11 @@ import dpctl.tensor as dpt import numpy +from numpy.core.numeric import normalize_axis_index import dpnp from dpnp.dpnp_algo import * from dpnp.dpnp_array import dpnp_array -from dpnp.dpnp_iface_arraycreation import array from dpnp.dpnp_utils import * __all__ = [ @@ -58,12 +58,16 @@ "concatenate", "copyto", "expand_dims", + "flip", + "fliplr", + "flipud", "hstack", "moveaxis", "ravel", "repeat", "reshape", "result_type", + "roll", "rollaxis", "shape", "squeeze", @@ -75,7 +79,7 @@ ] -def asfarray(x1, dtype=None): +def asfarray(a, dtype=None, *, device=None, usm_type=None, sycl_queue=None): """ Return an array converted to a float type. @@ -83,24 +87,37 @@ def asfarray(x1, dtype=None): Notes ----- - This function works exactly the same as :obj:`dpnp.array`. - If dtype is `None`, `bool` or one of the `int` dtypes, it is replaced with - the default floating type in DPNP depending on device capabilities. + If `dtype` is ``None``, :obj:`dpnp.bool` or one of the `int` dtypes, + it is replaced with the default floating type (:obj:`dpnp.float64` + if a device supports it, or :obj:`dpnp.float32` type otherwise). - """ + Returns + ------- + y : dpnp.ndarray + The input a as a float ndarray. + + Examples + -------- + >>> import dpnp as np + >>> np.asfarray([2, 3]) + array([2., 3.]) + >>> np.asfarray([2, 3], dtype=dpnp.float32) + array([2., 3.], dtype=float32) + >>> np.asfarray([2, 3], dtype=dpnp.int32) + array([2., 3.]) - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - if x1_desc: - if dtype is None or not numpy.issubdtype(dtype, dpnp.inexact): - dtype = dpnp.default_float_type(sycl_queue=x1.sycl_queue) + """ - # if type is the same then same object should be returned - if x1_desc.dtype == dtype: - return x1 + _sycl_queue = dpnp.get_normalized_queue_device( + a, sycl_queue=sycl_queue, device=device + ) - return array(x1, dtype=dtype) + if dtype is None or not numpy.issubdtype(dtype, dpnp.inexact): + dtype = dpnp.default_float_type(sycl_queue=_sycl_queue) - return call_origin(numpy.asfarray, x1, dtype) + return dpnp.asarray( + a, dtype=dtype, usm_type=usm_type, sycl_queue=_sycl_queue + ) def atleast_1d(*arys): @@ -112,13 +129,52 @@ def atleast_1d(*arys): For full documentation refer to :obj:`numpy.atleast_1d`. - Limitations - ----------- - Input arrays is supported as :obj:`dpnp.ndarray`. + Parameters + ---------- + arys : {dpnp_array, usm_ndarray} + One or more input arrays. + + Returns + ------- + out : dpnp.ndarray + An array, or list of arrays, each with ``a.ndim >= 1``. + Copies are made only if necessary. + + See Also + -------- + atleast_2d, atleast_3d + + Examples + -------- + >>> import dpnp as np + >>> np.atleast_1d(1.0) + array([1.]) + + >>> x = np.arange(9.0).reshape(3,3) + >>> np.atleast_1d(x) + array([[0., 1., 2.], + [3., 4., 5.], + [6., 7., 8.]]) + >>> np.atleast_1d(x) is x + True + + >>> np.atleast_1d(1, [3, 4]) + [array([1]), array([3, 4])] """ - return call_origin(numpy.atleast_1d, *arys) + res = [] + for ary in arys: + ary = dpnp.asanyarray(ary) + if ary.ndim == 0: + result = ary.reshape(1) + else: + result = ary + res.append(result) + if len(res) == 1: + return res[0] + else: + return res def atleast_2d(*arys): @@ -197,7 +253,7 @@ def atleast_3d(*arys): return call_origin(numpy.atleast_3d, *arys) -def broadcast_to(x, /, shape, subok=False): +def broadcast_to(array, /, shape, subok=False): """ Broadcast an array to a new shape. @@ -206,15 +262,15 @@ def broadcast_to(x, /, shape, subok=False): Returns ------- y : dpnp.ndarray - An array having a specified shape. Must have the same data type as `x`. + An array having a specified shape. Must have the same data type as `array`. Limitations ----------- - Parameter `x` is supported as either :class:`dpnp.ndarray` + Parameter `array` is supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. Parameter `subok` is supported with default value. Otherwise the function will be executed sequentially on CPU. - Input array data types of `x` is limited by supported DPNP :ref:`Data types`. + Input array data types of `array` is limited by supported DPNP :ref:`Data types`. Examples -------- @@ -229,15 +285,17 @@ def broadcast_to(x, /, shape, subok=False): if subok is not False: pass - elif isinstance(x, dpnp_array) or isinstance(x, dpt.usm_ndarray): - dpt_array = x.get_array() if isinstance(x, dpnp_array) else x + elif dpnp.is_supported_array_type(array): + dpt_array = dpnp.get_usm_ndarray(array) new_array = dpt.broadcast_to(dpt_array, shape) return dpnp_array._create_from_usm_ndarray(new_array) - return call_origin(numpy.broadcast_to, x, shape=shape, subok=subok) + return call_origin(numpy.broadcast_to, array, shape=shape, subok=subok) -def concatenate(arrays, /, *, axis=0, out=None, dtype=None, **kwargs): +def concatenate( + arrays, /, *, axis=0, out=None, dtype=None, casting="same_kind" +): """ Join a sequence of arrays along an existing axis. @@ -254,7 +312,6 @@ def concatenate(arrays, /, *, axis=0, out=None, dtype=None, **kwargs): or :class:`dpctl.tensor.usm_ndarray`. Otherwise ``TypeError`` exception will be raised. Parameters `out` and `dtype are supported with default value. - Keyword argument ``kwargs`` is currently unsupported. Otherwise the function will be executed sequentially on CPU. See Also @@ -288,12 +345,12 @@ def concatenate(arrays, /, *, axis=0, out=None, dtype=None, **kwargs): """ - if kwargs: - pass - elif out is not None: + if out is not None: pass elif dtype is not None: pass + elif casting != "same_kind": + pass else: usm_arrays = [dpnp.get_usm_ndarray(x) for x in arrays] usm_res = dpt.concat(usm_arrays, axis=axis) @@ -305,7 +362,7 @@ def concatenate(arrays, /, *, axis=0, out=None, dtype=None, **kwargs): axis=axis, out=out, dtype=dtype, - **kwargs, + casting=casting, ) @@ -388,7 +445,7 @@ def copyto(dst, src, casting="same_kind", where=True): dst_usm[mask_usm] = src_usm[mask_usm] -def expand_dims(x1, axis): +def expand_dims(a, axis): """ Expand the shape of an array. @@ -397,11 +454,35 @@ def expand_dims(x1, axis): For full documentation refer to :obj:`numpy.expand_dims`. + Returns + ------- + out : dpnp.ndarray + An array with the number of dimensions increased. + A view is returned whenever possible. + + Limitations + ----------- + Parameters `a` is supported either as :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. + Input array data types are limited by supported DPNP :ref:`Data types`. + Otherwise ``TypeError`` exception will be raised. + + Notes + ----- + If `a` has rank (i.e, number of dimensions) `N`, a valid `axis` must reside + in the closed-interval `[-N-1, N]`. + If provided a negative `axis`, the `axis` position at which to insert a + singleton dimension is computed as `N + axis + 1`. + Hence, if provided `-1`, the resolved axis position is `N` (i.e., + a singleton dimension must be appended to the input array `a`). + If provided `-N-1`, the resolved axis position is `0` (i.e., a + singleton dimension is prepended to the input array `a`). + See Also -------- :obj:`dpnp.squeeze` : The inverse operation, removing singleton dimensions :obj:`dpnp.reshape` : Insert, remove, and combine dimensions, and resize existing ones - :obj:`dpnp.indexing`, :obj:`dpnp.atleast_1d`, :obj:`dpnp.atleast_2d`, :obj:`dpnp.atleast_3d` + :obj:`dpnp.atleast_1d`, :obj:`dpnp.atleast_2d`, :obj:`dpnp.atleast_3d` Examples -------- @@ -446,33 +527,249 @@ def expand_dims(x1, axis): """ - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - if x1_desc: - return dpnp_expand_dims(x1_desc, axis).get_pyobj() + dpt_array = dpnp.get_usm_ndarray(a) + return dpnp_array._create_from_usm_ndarray( + dpt.expand_dims(dpt_array, axis=axis) + ) + + +def flip(m, axis=None): + """ + Reverse the order of elements in an array along the given axis. + + The shape of the array is preserved, but the elements are reordered. + + For full documentation refer to :obj:`numpy.flip`. + + Returns + ------- + out : dpnp.ndarray + A view of `m` with the entries of axis reversed. + + Limitations + ----------- + Parameters `m` is supported either as :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. + Input array data types are limited by supported DPNP :ref:`Data types`. + Otherwise ``TypeError`` exception will be raised. + + See Also + -------- + :obj:`dpnp.flipud` : Flip an array vertically (axis=0). + :obj:`dpnp.fliplr` : Flip an array horizontally (axis=1). + + Examples + -------- + >>> import dpnp as np + >>> A = np.arange(8).reshape((2, 2, 2)) + >>> A + array([[[0, 1], + [2, 3]], + [[4, 5], + [6, 7]]]) + >>> np.flip(A, 0) + array([[[4, 5], + [6, 7]], + [[0, 1], + [2, 3]]]) + >>> np.flip(A, 1) + array([[[2, 3], + [0, 1]], + [[6, 7], + [4, 5]]]) + >>> np.flip(A) + array([[[7, 6], + [5, 4]], + [[3, 2], + [1, 0]]]) + >>> np.flip(A, (0, 2)) + array([[[5, 4], + [7, 6]], + [[1, 0], + [3, 2]]]) + >>> A = np.random.randn(3, 4, 5) + >>> np.all(np.flip(A, 2) == A[:, :, ::-1, ...]) + array(True) + + """ + + m_usm = dpnp.get_usm_ndarray(m) + return dpnp_array._create_from_usm_ndarray(dpt.flip(m_usm, axis=axis)) + + +def fliplr(m): + """ + Reverse the order of elements along axis 1 (left/right). + + For a 2-D array, this flips the entries in each row in the left/right + direction. Columns are preserved, but appear in a different order than + before. + + For full documentation refer to :obj:`numpy.fliplr`. + + Returns + ------- + out : dpnp.ndarray + A view of `m` with the columns reversed. - return call_origin(numpy.expand_dims, x1, axis) + Limitations + ----------- + Parameters `m` is supported either as :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. + Input array data types are limited by supported DPNP :ref:`Data types`. + Otherwise ``TypeError`` exception will be raised. + See Also + -------- + :obj:`dpnp.flipud` : Flip an array vertically (axis=0). + :obj:`dpnp.flip` : Flip array in one or more dimensions. -def hstack(tup): + Examples + -------- + >>> import dpnp as np + >>> A = np.diag(np.array([1., 2., 3.])) + >>> A + array([[1., 0., 0.], + [0., 2., 0.], + [0., 0., 3.]]) + >>> np.fliplr(A) + array([[0., 0., 1.], + [0., 2., 0.], + [3., 0., 0.]]) + + >>> A = np.random.randn(2, 3, 5) + >>> np.all(np.fliplr(A) == A[:, ::-1, ...]) + array(True) + + """ + + if not dpnp.is_supported_array_type(m): + raise TypeError( + "An array must be any of supported type, but got {}".format(type(m)) + ) + + if m.ndim < 2: + raise ValueError(f"Input must be >= 2-d, but got {m.ndim}") + return m[:, ::-1] + + +def flipud(m): + """ + Reverse the order of elements along axis 0 (up/down). + + For a 2-D array, this flips the entries in each column in the up/down + direction. Rows are preserved, but appear in a different order than before. + + For full documentation refer to :obj:`numpy.flipud`. + + Returns + ------- + out : dpnp.ndarray + A view of `m` with the rows reversed. + + Limitations + ----------- + Parameters `m` is supported either as :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. + Input array data types are limited by supported DPNP :ref:`Data types`. + Otherwise ``TypeError`` exception will be raised. + + See Also + -------- + :obj:`dpnp.fliplr` : Flip array in the left/right direction. + :obj:`dpnp.flip` : Flip array in one or more dimensions. + + Examples + -------- + >>> import dpnp as np + >>> A = np.diag(np.array([1., 2., 3.])) + >>> A + array([[1., 0., 0.], + [0., 2., 0.], + [0., 0., 3.]]) + >>> np.flipud(A) + array([[0., 0., 3.], + [0., 2., 0.], + [1., 0., 0.]]) + + >>> A = np.random.randn(2, 3, 5) + >>> np.all(np.flipud(A) == A[::-1, ...]) + array(True) + + >>> np.flipud(np.array([1, 2])) + array([2, 1]) + + """ + + if not dpnp.is_supported_array_type(m): + raise TypeError( + "An array must be any of supported type, but got {}".format(type(m)) + ) + + if m.ndim < 1: + raise ValueError(f"Input must be >= 1-d, but got {m.ndim}") + return m[::-1, ...] + + +def hstack(tup, *, dtype=None, casting="same_kind"): """ Stack arrays in sequence horizontally (column wise). For full documentation refer to :obj:`numpy.hstack`. - """ + Returns + ------- + out : dpnp.ndarray + The stacked array which has one more dimension than the input arrays. - # TODO: - # `call_origin` cannot convert sequence of array to sequence of - # nparrays - tup_new = [] - for tp in tup: - tpx = dpnp.asnumpy(tp) if not isinstance(tp, numpy.ndarray) else tp - tup_new.append(tpx) + Limitations + ----------- + Each array in `tup` is supported as either :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. Otherwise ``TypeError`` exception + will be raised. + Parameters `dtype` and `casting` are supported with default value. + Otherwise the function will be executed sequentially on CPU. - return call_origin(numpy.hstack, tup_new) + See Also + -------- + :obj:`dpnp.concatenate` : Join a sequence of arrays along an existing axis. + :obj:`dpnp.stack` : Join a sequence of arrays along a new axis. + :obj:`dpnp.vstack` : Stack arrays in sequence vertically (row wise). + :obj:`dpnp.block` : Assemble an nd-array from nested lists of blocks. + :obj:`dpnp.split` : Split array into a list of multiple sub-arrays of equal size. + + Examples + -------- + >>> import dpnp as np + >>> a = np.array((1,2,3)) + >>> b = np.array((4,5,6)) + >>> np.hstack((a,b)) + array([1, 2, 3, 4, 5, 6]) + + >>> a = np.array([[1],[2],[3]]) + >>> b = np.array([[4],[5],[6]]) + >>> np.hstack((a,b)) + array([[1, 4], + [2, 5], + [3, 6]]) + + """ + + if not hasattr(tup, "__getitem__"): + raise TypeError( + "Arrays to stack must be passed as a sequence type such as list or tuple." + ) + arrs = dpnp.atleast_1d(*tup) + if not isinstance(arrs, list): + arrs = [arrs] + # As a special case, dimension 0 of 1-dimensional arrays is "horizontal" + if arrs and arrs[0].ndim == 1: + return dpnp.concatenate(arrs, axis=0, dtype=dtype, casting=casting) + else: + return dpnp.concatenate(arrs, axis=1, dtype=dtype, casting=casting) -def moveaxis(x, source, destination): +def moveaxis(a, source, destination): """ Move axes of an array to new positions. Other axes remain in their original order. @@ -483,14 +780,15 @@ def moveaxis(x, source, destination): out : dpnp.ndarray Array with moved axes. The returned array will have the same data and - the same USM allocation type as `x`. + the same USM allocation type as `a`. Limitations ----------- - Parameters `x` is supported as either :class:`dpnp.ndarray` - or :class:`dpctl.tensor.usm_ndarray`. - Otherwise the function will be executed sequentially on CPU. + Parameters `a` is supported as either :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. Otherwise ``TypeError`` exception + will be raised. Input array data types are limited by supported DPNP :ref:`Data types`. + Otherwise ``TypeError`` exception will be raised. See Also -------- @@ -508,16 +806,13 @@ def moveaxis(x, source, destination): """ - if isinstance(x, dpnp_array) or isinstance(x, dpt.usm_ndarray): - dpt_array = x.get_array() if isinstance(x, dpnp_array) else x - return dpnp_array._create_from_usm_ndarray( - dpt.moveaxis(dpt_array, source, destination) - ) - - return call_origin(numpy.moveaxis, x, source, destination) + dpt_array = dpnp.get_usm_ndarray(a) + return dpnp_array._create_from_usm_ndarray( + dpt.moveaxis(dpt_array, source, destination) + ) -def ravel(x1, order="C"): +def ravel(a, order="C"): """ Return a contiguous flattened array. @@ -539,14 +834,14 @@ def ravel(x1, order="C"): """ - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - if x1_desc: - return dpnp_flatten(x1_desc).get_pyobj() + a_desc = dpnp.get_dpnp_descriptor(a, copy_when_nondefault_queue=False) + if a_desc: + return dpnp_flatten(a_desc).get_pyobj() - return call_origin(numpy.ravel, x1, order=order) + return call_origin(numpy.ravel, a, order=order) -def repeat(x1, repeats, axis=None): +def repeat(a, repeats, axis=None): """ Repeat elements of an array. @@ -572,22 +867,22 @@ def repeat(x1, repeats, axis=None): """ - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - if x1_desc: + a_desc = dpnp.get_dpnp_descriptor(a, copy_when_nondefault_queue=False) + if a_desc: if axis is not None and axis != 0: pass - elif x1_desc.ndim >= 2: + elif a_desc.ndim >= 2: pass elif not dpnp.isscalar(repeats) and len(repeats) > 1: pass else: repeat_val = repeats if dpnp.isscalar(repeats) else repeats[0] - return dpnp_repeat(x1_desc, repeat_val, axis).get_pyobj() + return dpnp_repeat(a_desc, repeat_val, axis).get_pyobj() - return call_origin(numpy.repeat, x1, repeats, axis) + return call_origin(numpy.repeat, a, repeats, axis) -def reshape(x, /, newshape, order="C", copy=None): +def reshape(a, /, newshape, order="C", copy=None): """ Gives a new shape to an array without changing its data. @@ -595,7 +890,7 @@ def reshape(x, /, newshape, order="C", copy=None): Parameters ---------- - x : {dpnp_array, usm_ndarray} + a : {dpnp_array, usm_ndarray} Array to be reshaped. newshape : int or tuple of ints The new shape should be compatible with the original shape. If @@ -603,7 +898,7 @@ def reshape(x, /, newshape, order="C", copy=None): One shape dimension can be -1. In this case, the value is inferred from the length of the array and remaining dimensions. order : {'C', 'F'}, optional - Read the elements of `x` using this index order, and place the + Read the elements of `a` using this index order, and place the elements into the reshaped array using this index order. 'C' means to read / write the elements using C-like index order, with the last axis index changing fastest, back to the first @@ -614,10 +909,10 @@ def reshape(x, /, newshape, order="C", copy=None): the underlying array, and only refer to the order of indexing. copy : bool, optional Boolean indicating whether or not to copy the input array. - If ``True``, the result array will always be a copy of input `x`. + If ``True``, the result array will always be a copy of input `a`. If ``False``, the result array can never be a copy and a ValueError exception will be raised in case the copy is necessary. - If ``None``, the result array will reuse existing memory buffer of `x` + If ``None``, the result array will reuse existing memory buffer of `a` if possible and copy otherwise. Default: None. Returns @@ -652,14 +947,14 @@ def reshape(x, /, newshape, order="C", copy=None): """ if newshape is None: - newshape = x.shape + newshape = a.shape if order is None: order = "C" elif order not in "cfCF": raise ValueError(f"order must be one of 'C' or 'F' (got {order})") - usm_arr = dpnp.get_usm_ndarray(x) + usm_arr = dpnp.get_usm_ndarray(a) usm_arr = dpt.reshape(usm_arr, shape=newshape, order=order, copy=copy) return dpnp_array._create_from_usm_ndarray(usm_arr) @@ -709,18 +1004,81 @@ def result_type(*arrays_and_dtypes): return dpt.result_type(*usm_arrays_and_dtypes) -def rollaxis(x1, axis, start=0): +def roll(x, shift, axis=None): + """ + Roll the elements of an array by a number of positions along a given axis. + + Array elements that roll beyond the last position are re-introduced + at the first position. Array elements that roll beyond the first position + are re-introduced at the last position. + + For full documentation refer to :obj:`numpy.roll`. + + Returns + ------- + dpnp.ndarray + An array with the same data type as `x` + and whose elements, relative to `x`, are shifted. + + Limitations + ----------- + Parameter `x` is supported either as :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. Otherwise ``TypeError`` exception + will be raised. + Input array data types are limited by supported DPNP :ref:`Data types`. + + + See Also + -------- + :obj:`dpnp.moveaxis` : Move array axes to new positions. + :obj:`dpnp.rollaxis` : Roll the specified axis backwards + until it lies in a given position. + + Examples + -------- + >>> import dpnp as np + >>> x1 = np.arange(10) + >>> np.roll(x1, 2) + array([8, 9, 0, 1, 2, 3, 4, 5, 6, 7]) + + >>> np.roll(x1, -2) + array([2, 3, 4, 5, 6, 7, 8, 9, 0, 1]) + + >>> x2 = np.reshape(x1, (2, 5)) + >>> np.roll(x2, 1, axis=0) + array([[5, 6, 7, 8, 9], + [0, 1, 2, 3, 4]]) + + >>> np.roll(x2, (2, 1), axis=(1, 0)) + array([[8, 9, 5, 6, 7], + [3, 4, 0, 1, 2]]) + + """ + if axis is None: + return roll(x.reshape(-1), shift, 0).reshape(x.shape) + dpt_array = dpnp.get_usm_ndarray(x) + return dpnp_array._create_from_usm_ndarray( + dpt.roll(dpt_array, shift=shift, axis=axis) + ) + + +def rollaxis(x, axis, start=0): """ Roll the specified axis backwards, until it lies in a given position. For full documentation refer to :obj:`numpy.rollaxis`. + Returns + ------- + dpnp.ndarray + An array with the same data type as `x` where the specified axis + has been repositioned to the desired position. + Limitations ----------- - Input array is supported as :obj:`dpnp.ndarray`. - Parameter ``axis`` is supported as integer only. - Parameter ``start`` is limited by ``-a.ndim <= start <= a.ndim``. - Otherwise the function will be executed sequentially on CPU. + Parameter `x` is supported either as :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. Otherwise ``TypeError`` exception + will be raised. Input array data types are limited by supported DPNP :ref:`Data types`. See Also @@ -742,19 +1100,19 @@ def rollaxis(x1, axis, start=0): """ - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - if x1_desc: - if not isinstance(axis, int): - pass - elif start < -x1_desc.ndim or start > x1_desc.ndim: - pass - else: - start_norm = start + x1_desc.ndim if start < 0 else start - destination = start_norm - 1 if start_norm > axis else start_norm - - return dpnp.moveaxis(x1_desc.get_pyobj(), axis, destination) - - return call_origin(numpy.rollaxis, x1, axis, start) + n = x.ndim + axis = normalize_axis_index(axis, n) + if start < 0: + start += n + msg = "'%s' arg requires %d <= %s < %d, but %d was passed in" + if not (0 <= start < n + 1): + raise ValueError(msg % ("start", -n, "start", n + 1, start)) + if axis < start: + start -= 1 + if axis == start: + return x + dpt_array = dpnp.get_usm_ndarray(x) + return dpnp.moveaxis(dpt_array, source=axis, destination=start) def shape(a): @@ -800,9 +1158,9 @@ def shape(a): return numpy.shape(a) -def squeeze(x, /, axis=None): +def squeeze(a, /, axis=None): """ - Removes singleton dimensions (axes) from array `x`. + Removes singleton dimensions (axes) from array `a`. For full documentation refer to :obj:`numpy.squeeze`. @@ -814,13 +1172,14 @@ def squeeze(x, /, axis=None): dimensions of length 1 removed. Output has the same data type as the input, is allocated on the same device as the input and has the same USM allocation type as the input - array `x`. + array `a`. Limitations ----------- - Parameters `x` is supported as either :class:`dpnp.ndarray` + Parameters `a` is supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. - Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported DPNP :ref:`Data types`. + Otherwise ``TypeError`` exception will be raised. Examples -------- @@ -841,11 +1200,10 @@ def squeeze(x, /, axis=None): """ - if isinstance(x, dpnp_array) or isinstance(x, dpt.usm_ndarray): - dpt_array = x.get_array() if isinstance(x, dpnp_array) else x - return dpnp_array._create_from_usm_ndarray(dpt.squeeze(dpt_array, axis)) - - return call_origin(numpy.squeeze, x, axis) + dpt_array = dpnp.get_usm_ndarray(a) + return dpnp_array._create_from_usm_ndarray( + dpt.squeeze(dpt_array, axis=axis) + ) def stack(arrays, /, *, axis=0, out=None, dtype=None, **kwargs): @@ -864,7 +1222,7 @@ def stack(arrays, /, *, axis=0, out=None, dtype=None, **kwargs): Each array in `arrays` is supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. Otherwise ``TypeError`` exception will be raised. - Parameters `out` and `dtype are supported with default value. + Parameters `out` and `dtype` are supported with default value. Keyword argument ``kwargs`` is currently unsupported. Otherwise the function will be executed sequentially on CPU. @@ -921,50 +1279,57 @@ def stack(arrays, /, *, axis=0, out=None, dtype=None, **kwargs): ) -def swapaxes(x1, axis1, axis2): +def swapaxes(a, axis1, axis2): """ Interchange two axes of an array. For full documentation refer to :obj:`numpy.swapaxes`. + Returns + ------- + out : dpnp.ndarray + An array with with swapped axes. + A view is returned whenever possible. + Limitations ----------- - Input array is supported as :obj:`dpnp.ndarray`. - Otherwise the function will be executed sequentially on CPU. - Parameter ``axis1`` is limited by ``axis1 < x1.ndim``. - Parameter ``axis2`` is limited by ``axis2 < x1.ndim``. + Parameters `a` is supported either as :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. Input array data types are limited by supported DPNP :ref:`Data types`. + Otherwise ``TypeError`` exception will be raised. + + Notes + ----- + If `a` has rank (i.e., number of dimensions) `N`, + a valid `axis` must be in the half-open interval `[-N, N)`. Examples -------- >>> import dpnp as np >>> x = np.array([[1, 2, 3]]) - >>> out = np.swapaxes(x, 0, 1) - >>> out.shape - (3, 1) - >>> [i for i in out] - [1, 2, 3] + >>> np.swapaxes(x, 0, 1) + array([[1], + [2], + [3]]) + + >>> x = np.array([[[0,1],[2,3]],[[4,5],[6,7]]]) + >>> x + array([[[0, 1], + [2, 3]], + [[4, 5], + [6, 7]]]) + >>> np.swapaxes(x,0,2) + array([[[0, 4], + [2, 6]], + [[1, 5], + [3, 7]]]) """ - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - if x1_desc: - if axis1 >= x1_desc.ndim: - pass - elif axis2 >= x1_desc.ndim: - pass - else: - # 'do nothing' pattern for transpose() - input_permute = [i for i in range(x1.ndim)] - # swap axes - input_permute[axis1], input_permute[axis2] = ( - input_permute[axis2], - input_permute[axis1], - ) - - return transpose(x1_desc.get_pyobj(), axes=input_permute) - - return call_origin(numpy.swapaxes, x1, axis1, axis2) + dpt_array = dpnp.get_usm_ndarray(a) + return dpnp_array._create_from_usm_ndarray( + dpt.swapaxes(dpt_array, axis1=axis1, axis2=axis2) + ) def transpose(a, axes=None): @@ -1030,7 +1395,7 @@ def transpose(a, axes=None): return array.transpose(*axes) -def unique(x1, **kwargs): +def unique(ar, **kwargs): """ Find the unique elements of an array. @@ -1046,7 +1411,7 @@ def unique(x1, **kwargs): """ - return call_origin(numpy.unique, x1, **kwargs) + return call_origin(numpy.unique, ar, **kwargs) def vstack(tup): diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 7552bcd21304..4fad31bf9ee3 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -58,8 +58,12 @@ dpnp_floor_divide, dpnp_multiply, dpnp_negative, + dpnp_power, + dpnp_proj, dpnp_remainder, + dpnp_round, dpnp_sign, + dpnp_signbit, dpnp_subtract, dpnp_trunc, ) @@ -100,9 +104,12 @@ "negative", "power", "prod", + "proj", "remainder", - "round_", + "rint", + "round", "sign", + "signbit", "subtract", "sum", "trapz", @@ -144,14 +151,14 @@ def absolute(x, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray An array containing the absolute value of each element in `x`. Limitations ----------- Parameters `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. Parameters `out`, `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -218,7 +225,7 @@ def add( Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray The sum of `x1` and `x2`, element-wise. Limitations @@ -226,7 +233,7 @@ def add( Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -268,45 +275,39 @@ def add( ) -def around(x1, decimals=0, out=None): +def around(x, /, decimals=0, out=None): """ - Evenly round to the given number of decimals. + Round an array to the given number of decimals. For full documentation refer to :obj:`numpy.around`. + Returns + ------- + out : dpnp.ndarray + The rounded value of elements of the array. + Limitations ----------- - Parameters `x1` is supported as :class:`dpnp.ndarray`. - Parameters `decimals` and `out` are supported with their default values. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `decimals` is supported with its default value. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. - Examples + See Also -------- - >>> import dpnp as np - >>> np.around([0.37, 1.64]) - array([0., 2.]) - >>> np.around([0.37, 1.64], decimals=1) - array([0.4, 1.6]) - >>> np.around([.5, 1.5, 2.5, 3.5, 4.5]) # rounds to nearest even value - array([0., 2., 2., 4., 4.]) - >>> np.around([1,2,3,11], decimals=1) # ndarray of ints is returned - array([ 1, 2, 3, 11]) - >>> np.around([1,2,3,11], decimals=-1) - array([ 0, 0, 0, 10]) + :obj:`dpnp.round` : Equivalent function; see for details. + :obj:`dpnp.ndarray.round` : Equivalent function. + :obj:`dpnp.rint` : Round elements of the array to the nearest integer. + :obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise. + :obj:`dpnp.floor` : Return the floor of the input, element-wise. + :obj:`dpnp.trunc` : Return the truncated value of the input, element-wise. + Notes + ----- + This function works the same as :obj:`dpnp.round`. """ - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - if x1_desc: - if out is not None: - pass - elif decimals != 0: - pass - else: - return dpnp_around(x1_desc, decimals).get_pyobj() - - return call_origin(numpy.around, x1, decimals=decimals, out=out) + return round(x, decimals, out) def ceil( @@ -334,7 +335,7 @@ def ceil( ----------- Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. Parameters `where`, `dtype`, and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by real-value data types. @@ -442,58 +443,96 @@ def convolve(a, v, mode="full"): return call_origin(numpy.convolve, a=a, v=v, mode=mode) -def copysign(x1, x2, dtype=None, out=None, where=True, **kwargs): +def copysign( + x1, x2, /, out=None, *, where=True, dtype=None, subok=True, **kwargs +): """ - Change the sign of x1 to that of x2, element-wise. + Change the sign of `x1` to that of `x2`, element-wise. For full documentation refer to :obj:`numpy.copysign`. + Returns + ------- + out : dpnp.ndarray + The values of `x1` with the sign of `x2`. + Limitations ----------- - Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar. - Parameters `dtype`, `out` and `where` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. Examples -------- >>> import dpnp as np - >>> result = np.copysign(np.array([1, -2, 6, -9]), np.array([-1, -1, 1, 1])) - >>> [x for x in result] - [-1.0, -2.0, 6.0, 9.0] + >>> np.copysign(np.array(1.3), np.array(-1)) + array(-1.3) + >>> 1 / np.copysign(np.array(0), 1) + array(inf) + >>> 1 / np.copysign(np.array(0), -1) + array(-inf) + + >>> x = np.array([-1, 0, 1]) + >>> np.copysign(x, -1.1) + array([-1., -0., -1.]) + >>> np.copysign(x, np.arange(3) - 1) + array([-1., 0., 1.]) """ - x1_is_scalar = dpnp.isscalar(x1) - x2_is_scalar = dpnp.isscalar(x2) - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False - ) - x2_desc = dpnp.get_dpnp_descriptor( - x2, copy_when_strides=False, copy_when_nondefault_queue=False - ) + if kwargs: + pass + elif where is not True: + pass + elif dtype is not None: + pass + elif subok is not True: + pass + elif dpnp.isscalar(x1) and dpnp.isscalar(x2): + # at least either x1 or x2 has to be an array + pass + else: + # get USM type and queue to copy scalar from the host memory into a USM allocation + usm_type, queue = ( + get_usm_allocations([x1, x2]) + if dpnp.isscalar(x1) or dpnp.isscalar(x2) + else (None, None) + ) + + x1_desc = dpnp.get_dpnp_descriptor( + x1, + copy_when_strides=False, + copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, + alloc_queue=queue, + ) + x2_desc = dpnp.get_dpnp_descriptor( + x2, + copy_when_strides=False, + copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, + alloc_queue=queue, + ) + if x1_desc and x2_desc: + if out is not None: + if not dpnp.is_supported_array_type(out): + raise TypeError( + "return array must be of supported array type" + ) + out_desc = ( + dpnp.get_dpnp_descriptor( + out, copy_when_nondefault_queue=False + ) + or None + ) + else: + out_desc = None - if x1_desc and x2_desc and not kwargs: - if not x1_desc and not x1_is_scalar: - pass - elif not x2_desc and not x2_is_scalar: - pass - elif x1_is_scalar and x2_is_scalar: - pass - elif x1_desc and x1_desc.ndim == 0: - pass - elif x2_desc and x2_desc.ndim == 0: - pass - elif dtype is not None: - pass - elif out is not None: - pass - elif not where: - pass - else: return dpnp_copysign( - x1_desc, x2_desc, dtype=dtype, out=out, where=where + x1_desc, x2_desc, dtype=dtype, out=out_desc, where=where ).get_pyobj() return call_origin( @@ -510,7 +549,7 @@ def cross(x1, x2, axisa=-1, axisb=-1, axisc=-1, axis=None): Limitations ----------- Parameters `x1` and `x2` are supported as :class:`dpnp.ndarray`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Sizes of input arrays are limited by `x1.size == 3 and x2.size == 3`. Shapes of input arrays are limited by `x1.shape == (3,) and x2.shape == (3,)`. Otherwise the function will be executed sequentially on CPU. @@ -558,7 +597,7 @@ def cumprod(x1, **kwargs): Limitations ----------- Parameter `x` is supported as :class:`dpnp.ndarray`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -592,7 +631,7 @@ def cumsum(x1, **kwargs): Limitations ----------- Parameter `x` is supported as :obj:`dpnp.ndarray`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -671,7 +710,7 @@ def divide( Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray The quotient `x1/x2`, element-wise. Limitations @@ -773,7 +812,7 @@ def fabs(x1, **kwargs): Limitations ----------- Parameter `x1` is supported as :class:`dpnp.ndarray`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -824,7 +863,7 @@ def floor( ----------- Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. Parameters `where`, `dtype`, and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by real-value data types. @@ -877,12 +916,17 @@ def floor_divide( For full documentation refer to :obj:`numpy.floor_divide`. + Returns + ------- + out : dpnp.ndarray + The floordivide of each element of `x`. + Limitations ----------- Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -901,6 +945,14 @@ def floor_divide( >>> np.floor_divide(np.array([1., 2., 3., 4.]), 2.5) array([ 0., 0., 1., 1.]) + + The ``//`` operator can be used as a shorthand for ``floor_divide`` on + :class:`dpnp.ndarray`. + + >>> x1 = np.array([1., 2., 3., 4.]) + >>> x1 // 2.5 + array([0., 0., 1., 1.]) + """ return check_nd_call_func( @@ -959,17 +1011,23 @@ def fmin(*args, **kwargs): return dpnp.minimum(*args, **kwargs) -def fmod(x1, x2, dtype=None, out=None, where=True, **kwargs): +def fmod(x1, x2, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): """ - Calculate the element-wise remainder of division. + Returns the element-wise remainder of division. For full documentation refer to :obj:`numpy.fmod`. + Returns + ------- + out : dpnp.ndarray + The remainder of the division of `x1` by `x2`. + Limitations ----------- - Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar. - Parameters `dtype`, `out` and `where` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -981,48 +1039,80 @@ def fmod(x1, x2, dtype=None, out=None, where=True, **kwargs): Examples -------- >>> import dpnp as np - >>> a = np.array([2, -3, 4, 5, -4.5]) - >>> b = np.array([2, 2, 2, 2, 2]) - >>> result = np.fmod(a, b) - >>> [x for x in result] - [0.0, -1.0, 0.0, 1.0, -0.5] + >>> a = np.array([-3, -2, -1, 1, 2, 3]) + >>> np.fmod(a, 2) + array([-1, 0, -1, 1, 0, 1]) + >>> np.remainder(a, 2) + array([1, 0, 1, 1, 0, 1]) + + >>> a = np.array([5, 3]) + >>> b = np.array([2, 2.]) + >>> np.fmod(a, b) + array([1., 1.]) + + >>> a = np.arange(-3, 3).reshape(3, 2) + >>> a + array([[-3, -2], + [-1, 0], + [ 1, 2]]) + >>> b = np.array([2, 2]) + >>> np.fmod(a, b) + array([[-1, 0], + [-1, 0], + [ 1, 0]]) """ - x1_is_scalar = dpnp.isscalar(x1) - x2_is_scalar = dpnp.isscalar(x2) - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False - ) - x2_desc = dpnp.get_dpnp_descriptor( - x2, copy_when_strides=False, copy_when_nondefault_queue=False - ) + if kwargs: + pass + elif where is not True: + pass + elif dtype is not None: + pass + elif subok is not True: + pass + elif dpnp.isscalar(x1) and dpnp.isscalar(x2): + # at least either x1 or x2 has to be an array + pass + else: + # get USM type and queue to copy scalar from the host memory into a USM allocation + usm_type, queue = ( + get_usm_allocations([x1, x2]) + if dpnp.isscalar(x1) or dpnp.isscalar(x2) + else (None, None) + ) + + x1_desc = dpnp.get_dpnp_descriptor( + x1, + copy_when_strides=False, + copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, + alloc_queue=queue, + ) + x2_desc = dpnp.get_dpnp_descriptor( + x2, + copy_when_strides=False, + copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, + alloc_queue=queue, + ) + if x1_desc and x2_desc: + if out is not None: + if not dpnp.is_supported_array_type(out): + raise TypeError( + "return array must be of supported array type" + ) + out_desc = ( + dpnp.get_dpnp_descriptor( + out, copy_when_nondefault_queue=False + ) + or None + ) + else: + out_desc = None - if x1_desc and x2_desc and not kwargs: - if not x1_desc and not x1_is_scalar: - pass - elif not x2_desc and not x2_is_scalar: - pass - elif x1_is_scalar and x2_is_scalar: - pass - elif x1_desc and x1_desc.ndim == 0: - pass - elif x2_desc and x2_desc.ndim == 0: - pass - elif dtype is not None: - pass - elif out is not None: - pass - elif not where: - pass - else: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) return dpnp_fmod( - x1_desc, x2_desc, dtype, out_desc, where + x1_desc, x2_desc, dtype=dtype, out=out_desc, where=where ).get_pyobj() return call_origin( @@ -1040,7 +1130,7 @@ def gradient(x1, *varargs, **kwargs): ----------- Parameter `y1` is supported as :class:`dpnp.ndarray`. Argument `varargs[0]` is supported as `int`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -1072,129 +1162,225 @@ def gradient(x1, *varargs, **kwargs): return call_origin(numpy.gradient, x1, *varargs, **kwargs) -def maximum(x1, x2, dtype=None, out=None, where=True, **kwargs): +def maximum( + x1, x2, /, out=None, *, where=True, dtype=None, subok=True, **kwargs +): """ Element-wise maximum of array elements. For full documentation refer to :obj:`numpy.maximum`. + Returns + ------- + out : dpnp.ndarray + The maximum of `x1` and `x2`, element-wise. + Limitations ----------- - Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar. - Parameters `dtype`, `out` and `where` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. See Also -------- - :obj:`dpnp.fmax` : Element-wise maximum of array elements. - :obj:`dpnp.fmin` : Element-wise minimum of array elements. - :obj:`dpnp.fmod` : Calculate the element-wise remainder of division. + :obj:`dpnp.minimum` : Element-wise minimum of two arrays, propagates NaNs. + :obj:`dpnp.fmax` : Element-wise maximum of two arrays, ignores NaNs. + :obj:`dpnp.amax` : The maximum value of an array along a given axis, propagates NaNs. + :obj:`dpnp.nanmax` : The maximum value of an array along a given axis, ignores NaNs. + :obj:`dpnp.fmin` : Element-wise minimum of two arrays, ignores NaNs. + :obj:`dpnp.amix` : The minimum value of an array along a given axis, propagates NaNs. + :obj:`dpnp.nanmix` : The minimum value of an array along a given axis, ignores NaNs. Example ------- >>> import dpnp as np - >>> result = np.fmax(np.array([-2, 3, 4]), np.array([1, 5, 2])) - >>> [x for x in result] - [1, 5, 4] - - """ - - x1_is_scalar = dpnp.isscalar(x1) - x2_is_scalar = dpnp.isscalar(x2) - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False - ) - x2_desc = dpnp.get_dpnp_descriptor( - x2, copy_when_strides=False, copy_when_nondefault_queue=False - ) + >>> x1 = np.array([2, 3, 4]) + >>> x2 = np.array([1, 5, 2]) + >>> np.maximum(x1, x2) + array([2, 5, 4]) - if x1_desc and x2_desc and not kwargs: - if not x1_desc and not x1_is_scalar: - pass - elif not x2_desc and not x2_is_scalar: - pass - elif x1_is_scalar and x2_is_scalar: - pass - elif x1_desc and x1_desc.ndim == 0: - pass - elif x2_desc and x2_desc.ndim == 0: - pass - elif dtype is not None: - pass - elif out is not None: - pass - elif not where: - pass - else: - return dpnp_maximum( - x1_desc, x2_desc, dtype=dtype, out=out, where=where - ).get_pyobj() + >>> x1 = np.eye(2) + >>> x2 = np.array([0.5, 2]) + >>> np.maximum(x1, x2) # broadcasting + array([[1. , 2. ], + [0.5, 2. ]]) - return call_origin( - numpy.maximum, x1, x2, dtype=dtype, out=out, where=where, **kwargs - ) + >>> x1 = np.array([np.nan, 0, np.nan]) + >>> x2 = np.array([0, np.nan, np.nan]) + >>> np.maximum(x1, x2) + array([ 0., 0., nan]) + >>> np.maximum(np.array(np.Inf), 1) + array(inf) -def minimum(x1, x2, dtype=None, out=None, where=True, **kwargs): """ - Element-wise minimum of array elements. - - For full documentation refer to :obj:`numpy.minimum`. - Limitations - ----------- - Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar. - Parameters `dtype`, `out` and `where` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. - Otherwise the function will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + if kwargs: + pass + elif where is not True: + pass + elif dtype is not None: + pass + elif subok is not True: + pass + elif dpnp.isscalar(x1) and dpnp.isscalar(x2): + # at least either x1 or x2 has to be an array + pass + else: + # get USM type and queue to copy scalar from the host memory into a USM allocation + usm_type, queue = ( + get_usm_allocations([x1, x2]) + if dpnp.isscalar(x1) or dpnp.isscalar(x2) + else (None, None) + ) - See Also - -------- - :obj:`dpnp.fmax` : Element-wise maximum of array elements. - :obj:`dpnp.fmin` : Element-wise minimum of array elements. - :obj:`dpnp.fmod` : Calculate the element-wise remainder of division. + x1_desc = dpnp.get_dpnp_descriptor( + x1, + copy_when_strides=False, + copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, + alloc_queue=queue, + ) + x2_desc = dpnp.get_dpnp_descriptor( + x2, + copy_when_strides=False, + copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, + alloc_queue=queue, + ) + if x1_desc and x2_desc: + if out is not None: + if not dpnp.is_supported_array_type(out): + raise TypeError( + "return array must be of supported array type" + ) + out_desc = ( + dpnp.get_dpnp_descriptor( + out, copy_when_nondefault_queue=False + ) + or None + ) + else: + out_desc = None + + return dpnp_maximum( + x1_desc, x2_desc, dtype=dtype, out=out_desc, where=where + ).get_pyobj() + + return call_origin( + numpy.maximum, x1, x2, dtype=dtype, out=out, where=where, **kwargs + ) + + +def minimum( + x1, x2, /, out=None, *, where=True, dtype=None, subok=True, **kwargs +): + """ + Element-wise minimum of array elements. + + For full documentation refer to :obj:`numpy.minimum`. + + Returns + ------- + out : dpnp.ndarray + The minimum of `x1` and `x2`, element-wise. + + Limitations + ----------- + Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported DPNP :ref:`Data types`. + + See Also + -------- + :obj:`dpnp.maximum` : Element-wise maximum of two arrays, propagates NaNs. + :obj:`dpnp.fmin` : Element-wise minimum of two arrays, ignores NaNs. + :obj:`dpnp.amin` : The minimum value of an array along a given axis, propagates NaNs. + :obj:`dpnp.nanmin` : The minimum value of an array along a given axis, ignores NaNs. + :obj:`dpnp.fmax` : Element-wise maximum of two arrays, ignores NaNs. + :obj:`dpnp.amax` : The maximum value of an array along a given axis, propagates NaNs. + :obj:`dpnp.nanmax` : The maximum value of an array along a given axis, ignores NaNs. Example ------- >>> import dpnp as np - >>> result = np.fmin(np.array([-2, 3, 4]), np.array([1, 5, 2])) - >>> [x for x in result] - [-2, 3, 2] + >>> x1 = np.array([2, 3, 4]) + >>> x2 = np.array([1, 5, 2]) + >>> np.minimum(x1, x2) + array([1, 3, 2]) + + >>> x1 = np.eye(2) + >>> x2 = np.array([0.5, 2]) + >>> np.minimum(x1, x2) # broadcasting + array([[0.5, 0. ], + [0. , 1. ]] + + >>> x1 = np.array([np.nan, 0, np.nan]) + >>> x2 = np.array([0, np.nan, np.nan]) + >>> np.minimum(x1, x2) + array([ 0., 0., nan]) + + >>> np.minimum(np.array(-np.Inf), 1) + array(-inf) """ - x1_is_scalar = dpnp.isscalar(x1) - x2_is_scalar = dpnp.isscalar(x2) - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False - ) - x2_desc = dpnp.get_dpnp_descriptor( - x2, copy_when_strides=False, copy_when_nondefault_queue=False - ) + if kwargs: + pass + elif where is not True: + pass + elif dtype is not None: + pass + elif subok is not True: + pass + elif dpnp.isscalar(x1) and dpnp.isscalar(x2): + # at least either x1 or x2 has to be an array + pass + else: + # get USM type and queue to copy scalar from the host memory into a USM allocation + usm_type, queue = ( + get_usm_allocations([x1, x2]) + if dpnp.isscalar(x1) or dpnp.isscalar(x2) + else (None, None) + ) + + x1_desc = dpnp.get_dpnp_descriptor( + x1, + copy_when_strides=False, + copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, + alloc_queue=queue, + ) + x2_desc = dpnp.get_dpnp_descriptor( + x2, + copy_when_strides=False, + copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, + alloc_queue=queue, + ) + if x1_desc and x2_desc: + if out is not None: + if not dpnp.is_supported_array_type(out): + raise TypeError( + "return array must be of supported array type" + ) + out_desc = ( + dpnp.get_dpnp_descriptor( + out, copy_when_nondefault_queue=False + ) + or None + ) + else: + out_desc = None - if x1_desc and x2_desc and not kwargs: - if not x1_desc and not x1_is_scalar: - pass - elif not x2_desc and not x2_is_scalar: - pass - elif x1_is_scalar and x2_is_scalar: - pass - elif x1_desc and x1_desc.ndim == 0: - pass - elif x2_desc and x2_desc.ndim == 0: - pass - elif dtype is not None: - pass - elif out is not None: - pass - elif not where: - pass - else: return dpnp_minimum( - x1_desc, x2_desc, dtype=dtype, out=out, where=where + x1_desc, x2_desc, dtype=dtype, out=out_desc, where=where ).get_pyobj() return call_origin( @@ -1219,12 +1405,17 @@ def mod( For full documentation refer to :obj:`numpy.mod`. + Returns + ------- + out : dpnp.ndarray + The element-wise remainder of the quotient `floor_divide(x1, x2)`. + Limitations ----------- Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -1261,7 +1452,7 @@ def modf(x1, **kwargs): Limitations ----------- Parameter `x` is supported as :obj:`dpnp.ndarray`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -1302,7 +1493,7 @@ def multiply( Returns ------- - y : {dpnp.ndarray, scalar} + out : {dpnp.ndarray, scalar} The product of `x1` and `x2`, element-wise. Limitations @@ -1310,7 +1501,7 @@ def multiply( Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -1360,7 +1551,7 @@ def nancumprod(x1, **kwargs): Limitations ----------- Parameter `x` is supported as :class:`dpnp.ndarray`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -1397,7 +1588,7 @@ def nancumsum(x1, **kwargs): Limitations ----------- Parameter `x` is supported as :class:`dpnp.ndarray`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -1435,7 +1626,7 @@ def nanprod(x1, **kwargs): Limitations ----------- Parameter `x1` is supported as :obj:`dpnp.ndarray`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -1465,7 +1656,7 @@ def nansum(x1, **kwargs): Limitations ----------- Parameter `x1` is supported as :class:`dpnp.ndarray`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -1507,13 +1698,13 @@ def negative( Returns ------- out : dpnp.ndarray - The numerical negative of each element of `x`. + The numerical negative of each element of `x`. Limitations ----------- Parameters `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -1548,7 +1739,18 @@ def negative( ) -def power(x1, x2, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): +def power( + x1, + x2, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ First array elements raised to powers from second array, element-wise. @@ -1559,7 +1761,7 @@ def power(x1, x2, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray The bases in `x1` raised to the exponents in `x2`. Limitations @@ -1567,7 +1769,7 @@ def power(x1, x2, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. Parameters `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -1581,68 +1783,50 @@ def power(x1, x2, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): Example ------- >>> import dpnp as dp - >>> a = dp.array([1, 2, 3, 4, 5]) - >>> b = dp.array([2, 2, 2, 2, 2]) - >>> result = dp.power(a, b) - >>> [x for x in result] - [1, 4, 9, 16, 25] + >>> a = dp.arange(6) + >>> dp.power(a, 3) + array([ 0, 1, 8, 27, 64, 125]) - """ + Raise the bases to different exponents. - if kwargs: - pass - elif where is not True: - pass - elif dtype is not None: - pass - elif subok is not True: - pass - elif dpnp.isscalar(x1) and dpnp.isscalar(x2): - # at least either x1 or x2 has to be an array - pass - else: - # get USM type and queue to copy scalar from the host memory into a USM allocation - usm_type, queue = ( - get_usm_allocations([x1, x2]) - if dpnp.isscalar(x1) or dpnp.isscalar(x2) - else (None, None) - ) + >>> b = dp.array([1.0, 2.0, 3.0, 3.0, 2.0, 1.0]) + >>> dp.power(a, b) + array([ 0., 1., 8., 27., 16., 5.]) - x1_desc = dpnp.get_dpnp_descriptor( - x1, - copy_when_strides=False, - copy_when_nondefault_queue=False, - alloc_usm_type=usm_type, - alloc_queue=queue, - ) - x2_desc = dpnp.get_dpnp_descriptor( - x2, - copy_when_strides=False, - copy_when_nondefault_queue=False, - alloc_usm_type=usm_type, - alloc_queue=queue, - ) - if x1_desc and x2_desc: - if out is not None: - if not isinstance(out, (dpnp.ndarray, dpt.usm_ndarray)): - raise TypeError( - "return array must be of supported array type" - ) - out_desc = ( - dpnp.get_dpnp_descriptor( - out, copy_when_nondefault_queue=False - ) - or None - ) - else: - out_desc = None + The effect of broadcasting. - return dpnp_power( - x1_desc, x2_desc, dtype=dtype, out=out_desc, where=where - ).get_pyobj() + >>> c = dp.array([[1, 2, 3, 3, 2, 1], [1, 2, 3, 3, 2, 1]]) + >>> dp.power(a, c) + array([[ 0, 1, 8, 27, 16, 5], + [ 0, 1, 8, 27, 16, 5]]) - return call_origin( - numpy.power, x1, x2, dtype=dtype, out=out, where=where, **kwargs + The ``**`` operator can be used as a shorthand for ``power`` on + :class:`dpnp.ndarray`. + + >>> b = dp.array([1, 2, 3, 3, 2, 1]) + >>> a = dp.arange(6) + >>> a ** b + array([ 0, 1, 8, 27, 16, 5]) + + Negative values raised to a non-integral value will result in ``nan``. + + >>> d = dp.array([-1.0, -4.0]) + >>> dp.power(d, 1.5) + array([nan, nan]) + + """ + + return check_nd_call_func( + numpy.power, + dpnp_power, + x1, + x2, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) @@ -1706,6 +1890,123 @@ def prod( ) +def proj( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): + """ + Returns the projection of a number onto the Riemann sphere. + + For all infinite complex numbers (including the cases where one component is infinite and the other is `NaN`), + the function returns `(inf, 0.0)` or `(inf, -0.0)`. + For finite complex numbers, the input is returned. + All real-valued numbers are treated as complex numbers with positive zero imaginary part. + + Returns + ------- + out : dpnp.ndarray + The projection of each element of `x`. + + Limitations + ----------- + Parameters `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Input array data types are limited by supported DPNP :ref:`Data types`. + + See Also + -------- + :obj:`dpnp.abs` : Returns the magnitude of a complex number, element-wise. + :obj:`dpnp.conj` : Return the complex conjugate, element-wise. + + Examples + -------- + >>> import dpnp as np + >>> np.proj(np.array([1, -2.3, 2.1-1.7j])) + array([ 1. +0.j, -2.3+0.j, 2.1-1.7.j]) + + >>> np.proj(np.array([complex(1,np.inf), complex(1,-np.inf), complex(np.inf,-1),])) + array([inf+0.j, inf-0.j, inf-0.j]) + """ + + return check_nd_call_func( + None, + dpnp_proj, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, + ) + + +def rint( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): + """ + Round elements of the array to the nearest integer. + + For full documentation refer to :obj:`numpy.rint`. + + Returns + ------- + out : dpnp.ndarray + The rounded value of elements of the array to the nearest integer. + + Limitations + ----------- + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported DPNP :ref:`Data types`. + + See Also + -------- + :obj:`dpnp.round` : Evenly round to the given number of decimals. + :obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise. + :obj:`dpnp.floor` : Return the floor of the input, element-wise. + :obj:`dpnp.trunc` : Return the truncated value of the input, element-wise. + + Examples + -------- + >>> import dpnp as np + >>> a = np.array([-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]) + >>> np.rint(a) + array([-2., -2., -0., 0., 2., 2., 2.]) + + """ + + return check_nd_call_func( + numpy.rint, + dpnp_round, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, + ) + + def remainder( x1, x2, @@ -1723,6 +2024,11 @@ def remainder( For full documentation refer to :obj:`numpy.remainder`. + Returns + ------- + out : dpnp.ndarray + The element-wise remainder of the quotient `floor_divide(x1, x2)`. + Limitations ----------- Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` @@ -1771,19 +2077,57 @@ def remainder( ) -def round_(a, decimals=0, out=None): +def round(x, decimals=0, out=None): """ - Round an array to the given number of decimals. + Evenly round to the given number of decimals. + + For full documentation refer to :obj:`numpy.round`. - For full documentation refer to :obj:`numpy.round_`. + Returns + ------- + out : dpnp.ndarray + The rounded value of elements of the array. + + Limitations + ----------- + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `decimals` is supported with its default value. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported DPNP :ref:`Data types`. See Also -------- - :obj:`dpnp.around` : equivalent function; see for details. + :obj:`dpnp.around` : Equivalent function; see for details. + :obj:`dpnp.ndarray.round` : Equivalent function. + :obj:`dpnp.rint` : Round elements of the array to the nearest integer. + :obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise. + :obj:`dpnp.floor` : Return the floor of the input, element-wise. + :obj:`dpnp.trunc` : Return the truncated value of the input, element-wise. + + Examples + -------- + >>> import dpnp as np + >>> np.round(np.array([0.37, 1.64])) + array([0., 2.]) + >>> np.round(np.array([0.37, 1.64]), decimals=1) + array([0.4, 1.6]) + >>> np.round(np.array([.5, 1.5, 2.5, 3.5, 4.5])) # rounds to nearest even value + array([0., 2., 2., 4., 4.]) + >>> np.round(np.array([1,2,3,11]), decimals=1) # ndarray of ints is returned + array([ 1, 2, 3, 11]) + >>> np.round(np.array([1,2,3,11]), decimals=-1) + array([ 0, 0, 0, 10]) """ - return around(a, decimals, out) + if decimals != 0: + pass + elif dpnp.isscalar(x): + # input has to be an array + pass + else: + return dpnp_round(x, out=out) + return call_origin(numpy.round, x, decimals=decimals, out=out) def sign( @@ -1805,7 +2149,7 @@ def sign( Returns ------- out : dpnp.ndarray - The indication of the sign of each element of `x`. + The indication of the sign of each element of `x`. Limitations ----------- @@ -1816,6 +2160,10 @@ def sign( Input array data types are limited by supported DPNP :ref:`Data types`. However, if the input array data type is complex, the function will be executed sequentially on CPU. + See Also + -------- + :obj:`dpnp.signbit` : Returns element-wise `True` where signbit is set (less than zero). + Examples -------- >>> import dpnp as np @@ -1853,6 +2201,63 @@ def sign( ) +def signbit( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): + """ + Returns element-wise `True` where signbit is set (less than zero). + + For full documentation refer to :obj:`numpy.signbit`. + + Returns + ------- + out : dpnp.ndarray + A boolean array with indication of the sign of each element of `x`. + + Limitations + ----------- + Parameters `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported real-valued data types. + + See Also + -------- + :obj:`dpnp.sign` : Returns an element-wise indication of the sign of a number. + + Examples + -------- + >>> import dpnp as np + >>> np.signbit(np.array([-1.2])) + array([True]) + + >>> np.signbit(np.array([1, -2.3, 2.1])) + array([False, True, False]) + + """ + + return check_nd_call_func( + numpy.signbit, + dpnp_signbit, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, + ) + + def subtract( x1, x2, @@ -1872,7 +2277,7 @@ def subtract( Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray The difference of `x1` and `x2`, element-wise. Limitations @@ -1938,7 +2343,7 @@ def sum( Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray an array containing the sums. If the sum was computed over the entire array, a zero-dimensional array is returned. The returned array has the data type as described in the `dtype` parameter @@ -2042,7 +2447,7 @@ def trapz(y1, x1=None, dx=1.0, axis=-1): Limitations ----------- Parameters `y` and `x` are supported as :class:`dpnp.ndarray`. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -2107,7 +2512,6 @@ def true_divide(*args, **kwargs): ----- This function works the same as :obj:`dpnp.divide`. - """ return dpnp.divide(*args, **kwargs) @@ -2138,7 +2542,7 @@ def trunc( ----------- Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. Parameters `where`, `dtype`, and `subok` are supported with their default values. - Keyword arguments `kwargs` are currently unsupported. + Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by real-value data types. diff --git a/dpnp/dpnp_iface_trigonometric.py b/dpnp/dpnp_iface_trigonometric.py index dc91fa75edaf..3ed6de42e998 100644 --- a/dpnp/dpnp_iface_trigonometric.py +++ b/dpnp/dpnp_iface_trigonometric.py @@ -48,11 +48,22 @@ from .dpnp_algo.dpnp_elementwise_common import ( check_nd_call_func, + dpnp_acos, + dpnp_acosh, + dpnp_asin, + dpnp_asinh, + dpnp_atan, + dpnp_atan2, + dpnp_atanh, dpnp_cos, + dpnp_cosh, dpnp_log, dpnp_sin, + dpnp_sinh, dpnp_sqrt, dpnp_square, + dpnp_tan, + dpnp_tanh, ) __all__ = [ @@ -89,326 +100,458 @@ ] -def arccos(x1): +def arccos( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Trigonometric inverse cosine, element-wise. For full documentation refer to :obj:`numpy.arccos`. + Returns + ------- + out : dpnp.ndarray + The inverse cosine of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. See Also -------- - :obj:`dpnp.cos` : Cosine element-wise. + :obj:`dpnp.cos` : Trigonometric cosine, element-wise. :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. - :obj:`dpnp.arcsin` : Inverse sine, element-wise. + :obj:`dpnp.arcsin` : Trigonometric inverse sine, element-wise. + :obj:`dpnp.arccosh` : Hyperbolic inverse cosine, element-wise. Examples -------- >>> import dpnp as np >>> x = np.array([1, -1]) - >>> out = np.arccos(x) - >>> [i for i in out] - [0.0, 3.14159265] + >>> np.arccos(x) + array([0.0, 3.14159265]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arccos, + dpnp_acos, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_arccos(x1_desc).get_pyobj() - - return call_origin(numpy.arccos, x1, **kwargs) -def arccosh(x1): +def arccosh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ - Trigonometric inverse hyperbolic cosine, element-wise. + Inverse hyperbolic cosine, element-wise. For full documentation refer to :obj:`numpy.arccosh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic inverse cosine of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. See Also -------- :obj:`dpnp.cosh` : Hyperbolic cosine, element-wise. - :obj:`dpnp.arcsinh` : Inverse hyperbolic sine element-wise. - :obj:`dpnp.sinh` : Hyperbolic sine, element-wise. - :obj:`dpnp.arctanh` : Inverse hyperbolic tangent element-wise. - :obj:`dpnp.tanh` : Compute hyperbolic tangent element-wise. + :obj:`dpnp.arctanh` : Hyperbolic inverse tangent, element-wise. + :obj:`dpnp.arcsinh` : Hyperbolic inverse sine, element-wise. + :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. Examples -------- >>> import dpnp as np - >>> x = np.array([np.e, 10.0]) - >>> out = np.arccosh(x) - >>> [i for i in out] - [1.65745445, 2.99322285] + >>> x = np.array([1.0, np.e, 10.0]) + >>> np.arccosh(x) + array([0.0, 1.65745445, 2.99322285]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arccosh, + dpnp_acosh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_arccosh(x1_desc).get_pyobj() - return call_origin(numpy.arccosh, x1, **kwargs) - -def arcsin(x1, out=None, **kwargs): +def arcsin( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Inverse sine, element-wise. For full documentation refer to :obj:`numpy.arcsin`. + Returns + ------- + out : dpnp.ndarray + The inverse sine of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. See Also -------- :obj:`dpnp.sin` : Trigonometric sine, element-wise. - :obj:`dpnp.cos` : Cosine element-wise. - :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. - :obj:`dpnp.tan` : Compute tangent element-wise. :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. - :obj:`dpnp.arctan2` : Element-wise arc tangent of ``x1/x2`` - choosing the quadrant correctly. + :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. + :obj:`dpnp.arcsinh` : Hyperbolic inverse sine, element-wise. Examples -------- >>> import dpnp as np >>> x = np.array([0, 1, -1]) - >>> out = np.arcsin(x) - >>> [i for i in out] - [0.0, 1.5707963267948966, -1.5707963267948966] + >>> np.arcsin(x) + array([0.0, 1.5707963267948966, -1.5707963267948966]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arcsin, + dpnp_asin, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_arcsin(x1_desc, out_desc).get_pyobj() - - return call_origin(numpy.arcsin, x1, out=out, **kwargs) -def arcsinh(x1): +def arcsinh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Inverse hyperbolic sine, element-wise. For full documentation refer to :obj:`numpy.arcsinh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic inverse sine of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.sinh` : Hyperbolic sine, element-wise. + :obj:`dpnp.arctanh` : Hyperbolic inverse tangent, element-wise. + :obj:`dpnp.arccosh` : Hyperbolic inverse cosine, element-wise. + :obj:`dpnp.arcsin` : Trigonometric inverse sine, element-wise. + Examples -------- >>> import dpnp as np >>> x = np.array([np.e, 10.0]) - >>> out = np.arcsinh(x) - >>> [i for i in out] - [1.72538256, 2.99822295] + >>> np.arcsinh(x) + array([1.72538256, 2.99822295]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arcsinh, + dpnp_asinh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_arcsinh(x1_desc).get_pyobj() - return call_origin(numpy.arcsinh, x1, **kwargs) - -def arctan(x1, out=None, **kwargs): +def arctan( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Trigonometric inverse tangent, element-wise. For full documentation refer to :obj:`numpy.arctan`. + Returns + ------- + out : dpnp.ndarray + The inverse tangent of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. Keyword argument `kwargs` is currently unsupported. - Input array data types are limited by supported DPNP :ref:`Data types`. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported real-valued floating-point data type. See Also -------- - :obj:`dpnp.arctan2` : Element-wise arc tangent of ``x1/x2`` - choosing the quadrant correctly. + :obj:`dpnp.arctan2` : Element-wise arc tangent of `x1/x2` choosing the quadrant correctly. :obj:`dpnp.angle` : Argument of complex values. + :obj:`dpnp.tan` : Trigonometric tangent, element-wise. + :obj:`dpnp.arcsin` : Trigonometric inverse sine, element-wise. + :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. + :obj:`dpnp.arctanh` : Inverse hyperbolic tangent, element-wise. Examples -------- >>> import dpnp as np >>> x = np.array([0, 1]) - >>> out = np.arctan(x) - >>> [i for i in out] - [0.0, 0.78539816] + >>> np.arctan(x) + array([0.0, 0.78539816]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arctan, + dpnp_atan, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_arctan(x1_desc, out_desc).get_pyobj() - return call_origin(numpy.arctan, x1, out=out, **kwargs) - -def arctanh(x1): +def arctan2( + x1, + x2, + /, + out=None, + *, + where=True, + order="K", + dtype=None, + subok=True, + **kwargs, +): """ - Trigonometric hyperbolic inverse tangent, element-wise. + Element-wise arc tangent of `x1/x2` choosing the quadrant correctly. - For full documentation refer to :obj:`numpy.arctanh`. + For full documentation refer to :obj:`numpy.arctan2`. + + Returns + ------- + out : dpnp.ndarray + The inverse tangent of `x1/x2`, element-wise. Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword arguments `kwargs` are currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. + :obj:`dpnp.tan` : Compute tangent element-wise. + :obj:`dpnp.angle` : Return the angle of the complex argument. + :obj:`dpnp.arcsin` : Trigonometric inverse sine, element-wise. + :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. + :obj:`dpnp.arctanh` : Inverse hyperbolic tangent, element-wise. + Examples -------- >>> import dpnp as np - >>> x = np.array([0, -0.5]) - >>> out = np.arctanh(x) - >>> [i for i in out] - [0.0, -0.54930614] + >>> x1 = np.array([1., -1.]) + >>> x2 = np.array([0., 0.]) + >>> np.arctan2(x1, x2) + array([1.57079633, -1.57079633]) + + >>> x1 = np.array([0., 0., np.inf]) + >>> x2 = np.array([+0., -0., np.inf]) + >>> np.arctan2(x1, x2) + array([0.0 , 3.14159265, 0.78539816]) + + >>> x1 = np.array([-1, +1, +1, -1]) + >>> x2 = np.array([-1, -1, +1, +1]) + >>> np.arctan2(x1, x2) * 180 / np.pi + array([-135., -45., 45., 135.]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arctan2, + dpnp_atan2, + x1, + x2, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_arctanh(x1_desc).get_pyobj() - return call_origin(numpy.arctanh, x1, **kwargs) - -def cbrt(x1): +def arctanh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ - Return the cube-root of an array, element-wise. + Hyperbolic inverse tangent, element-wise. - For full documentation refer to :obj:`numpy.cbrt`. + For full documentation refer to :obj:`numpy.arctanh`. + + Returns + ------- + out : dpnp.ndarray + The hyperbolic inverse tangent of each element of `x`. Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.tanh` : Hyperbolic tangent, element-wise. + :obj:`dpnp.arcsinh` : Hyperbolic inverse sine, element-wise. + :obj:`dpnp.arccosh` : Hyperbolic inverse cosine, element-wise. + :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. + Examples -------- >>> import dpnp as np - >>> x = np.array([1, 8, 27]) - >>> out = np.cbrt(x) - >>> [i for i in out] - [1.0, 2.0, 3.0] + >>> x = np.array([0, -0.5]) + >>> np.arctanh(x) + array([0.0, -0.54930614]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arctanh, + dpnp_atanh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_cbrt(x1_desc).get_pyobj() - - return call_origin(numpy.cbrt, x1, **kwargs) -def arctan2(x1, x2, dtype=None, out=None, where=True, **kwargs): +def cbrt(x1): """ - Element-wise arc tangent of ``x1/x2`` choosing the quadrant correctly. + Return the cube-root of an array, element-wise. - For full documentation refer to :obj:`numpy.arctan2`. + For full documentation refer to :obj:`numpy.cbrt`. Limitations ----------- - Parameters `x1` and `x2` are supported as either :obj:`dpnp.ndarray` or scalar. - Parameters `dtype`, `out` and `where` are supported with their default values. - Keyword argument `kwargs` is currently unsupported. - Otherwise the function will be executed sequentially on CPU. + Input array is supported as :class:`dpnp.ndarray`. Input array data types are limited by supported DPNP :ref:`Data types`. - See Also - -------- - :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. - :obj:`dpnp.tan` : Compute tangent element-wise. - :obj:`dpnp.angle` : Return the angle of the complex argument. - Examples -------- >>> import dpnp as np - >>> x1 = np.array([1., -1.]) - >>> x2 = np.array([0., 0.]) - >>> out = np.arctan2(x1, x2) + >>> x = np.array([1, 8, 27]) + >>> out = np.cbrt(x) >>> [i for i in out] - [1.57079633, -1.57079633] + [1.0, 2.0, 3.0] """ - x1_is_scalar = dpnp.isscalar(x1) - x2_is_scalar = dpnp.isscalar(x2) x1_desc = dpnp.get_dpnp_descriptor( x1, copy_when_strides=False, copy_when_nondefault_queue=False ) - x2_desc = dpnp.get_dpnp_descriptor( - x2, copy_when_strides=False, copy_when_nondefault_queue=False - ) - - if x1_desc and x2_desc and not kwargs: - if not x1_desc and not x1_is_scalar: - pass - elif not x2_desc and not x2_is_scalar: - pass - elif x1_is_scalar and x2_is_scalar: - pass - elif x1_desc and x1_desc.ndim == 0: - pass - elif x2_desc and x2_desc.ndim == 0: - pass - elif dtype is not None: - pass - elif not where: - pass - else: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_arctan2( - x1_desc, x2_desc, dtype, out_desc, where - ).get_pyobj() + if x1_desc: + return dpnp_cbrt(x1_desc).get_pyobj() - return call_origin( - numpy.arctan2, x1, x2, dtype=dtype, out=out, where=where, **kwargs - ) + return call_origin(numpy.cbrt, x1, **kwargs) def cos( @@ -429,7 +572,7 @@ def cos( Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray The cosine of each element of `x`. Limitations @@ -440,6 +583,13 @@ def cos( Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. + :obj:`dpnp.sin` : Trigonometric sine, element-wise. + :obj:`dpnp.tan` : Trigonometric tangent, element-wise. + :obj:`dpnp.cosh` : Hyperbolic cosine, element-wise. + Examples -------- >>> import dpnp as np @@ -462,34 +612,63 @@ def cos( ) -def cosh(x1): +def cosh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ - Trigonometric hyperbolic cosine, element-wise. + Hyperbolic cosine, element-wise. For full documentation refer to :obj:`numpy.cosh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic cosine of each element of `x`. + Limitations ----------- - Input array is supported as :obj:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arccosh` : Hyperbolic inverse cosine, element-wise. + :obj:`dpnp.sinh` : Hyperbolic sine, element-wise. + :obj:`dpnp.tanh` : Hyperbolic tangent, element-wise. + :obj:`dpnp.cos` : Trigonometric cosine, element-wise. + + Examples -------- >>> import dpnp as np - >>> x = np.array([0]) - >>> out = np.cosh(x) - >>> [i for i in out] - [1.0] + >>> x = np.array([0, np.pi/2, np.pi]) + >>> np.cosh(x) + array([1.0, 2.5091786, 11.591953]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.cosh, + dpnp_cosh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_cosh(x1_desc).get_pyobj() - - return call_origin(numpy.cosh, x1, **kwargs) def deg2rad(x1): @@ -546,7 +725,7 @@ def degrees(x1): def exp(x1, out=None, **kwargs): """ - Trigonometric exponent, element-wise. + Calculate the exponential, element-wise. For full documentation refer to :obj:`numpy.exp`. @@ -586,7 +765,7 @@ def exp(x1, out=None, **kwargs): def exp2(x1): """ - Trigonometric exponent2, element-wise. + Calculate `2**p` for all `p` in the input array. For full documentation refer to :obj:`numpy.exp2`. @@ -621,7 +800,7 @@ def exp2(x1): def expm1(x1): """ - Trigonometric exponent minus 1, element-wise. + Return the exponential of the input array minus one, element-wise. For full documentation refer to :obj:`numpy.expm1`. @@ -651,16 +830,22 @@ def expm1(x1): return call_origin(numpy.expm1, x1) -def hypot(x1, x2, dtype=None, out=None, where=True, **kwargs): +def hypot(x1, x2, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): """ Given the "legs" of a right triangle, return its hypotenuse. For full documentation refer to :obj:`numpy.hypot`. + Returns + ------- + out : dpnp.ndarray + The hypotenuse of the triangle(s). + Limitations ----------- - Parameters `x1` and `x2` are supported as either :obj:`dpnp.ndarray` or scalar. - Parameters `dtype`, `out` and `where` are supported with their default values. + Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. + Parameters `where`, `dtype` and `subok` are supported with their default values. Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. @@ -668,48 +853,72 @@ def hypot(x1, x2, dtype=None, out=None, where=True, **kwargs): Examples -------- >>> import dpnp as np - >>> x1 = 3 * np.ones(3) - >>> x2 = 4 * np.ones(3) - >>> out = np.hypot(x1, x2) - >>> [i for i in out] - [5.0, 5.0, 5.0] + >>> x1 = 3 * np.ones((3, 3)) + >>> x2 = 4 * np.ones((3, 3)) + >>> np.hypot(x1, x2) + array([[5., 5., 5.], + [5., 5., 5.], + [5., 5., 5.]]) + + Example showing broadcast of scalar argument: + + >>> np.hypot(x1, 4) + array([[ 5., 5., 5.], + [ 5., 5., 5.], + [ 5., 5., 5.]]) """ - x1_is_scalar = dpnp.isscalar(x1) - x2_is_scalar = dpnp.isscalar(x2) - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False - ) - x2_desc = dpnp.get_dpnp_descriptor( - x2, copy_when_strides=False, copy_when_nondefault_queue=False - ) + if kwargs: + pass + elif where is not True: + pass + elif dtype is not None: + pass + elif subok is not True: + pass + elif dpnp.isscalar(x1) and dpnp.isscalar(x2): + # at least either x1 or x2 has to be an array + pass + else: + # get USM type and queue to copy scalar from the host memory into a USM allocation + usm_type, queue = ( + get_usm_allocations([x1, x2]) + if dpnp.isscalar(x1) or dpnp.isscalar(x2) + else (None, None) + ) + + x1_desc = dpnp.get_dpnp_descriptor( + x1, + copy_when_strides=False, + copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, + alloc_queue=queue, + ) + x2_desc = dpnp.get_dpnp_descriptor( + x2, + copy_when_strides=False, + copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, + alloc_queue=queue, + ) + if x1_desc and x2_desc: + if out is not None: + if not dpnp.is_supported_array_type(out): + raise TypeError( + "return array must be of supported array type" + ) + out_desc = ( + dpnp.get_dpnp_descriptor( + out, copy_when_nondefault_queue=False + ) + or None + ) + else: + out_desc = None - if x1_desc and x2_desc and not kwargs: - if not x1_desc and not x1_is_scalar: - pass - elif not x2_desc and not x2_is_scalar: - pass - elif x1_is_scalar and x2_is_scalar: - pass - elif x1_desc and x1_desc.ndim == 0: - pass - elif x2_desc and x2_desc.ndim == 0: - pass - elif dtype is not None: - pass - elif out is not None: - pass - elif not where: - pass - else: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) return dpnp_hypot( - x1_desc, x2_desc, dtype, out_desc, where + x1_desc, x2_desc, dtype=dtype, out=out_desc, where=where ).get_pyobj() return call_origin( @@ -738,7 +947,7 @@ def log( Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray The natural logarithm of `x`, element-wise. Limitations @@ -780,7 +989,7 @@ def log( def log10(x1): """ - Trigonometric logarithm, element-wise. + Return the base 10 logarithm of the input array, element-wise. For full documentation refer to :obj:`numpy.log10`. @@ -810,7 +1019,7 @@ def log10(x1): def log1p(x1): """ - Trigonometric logarithm, element-wise. + Return the natural logarithm of one plus the input array, element-wise. For full documentation refer to :obj:`numpy.log1p`. @@ -844,7 +1053,7 @@ def log1p(x1): def log2(x1): """ - Trigonometric logarithm, element-wise. + Return the base 2 logarithm of the input array, element-wise. For full documentation refer to :obj:`numpy.log2`. @@ -982,7 +1191,7 @@ def sin( Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray The sine of each element of `x`. Limitations @@ -995,9 +1204,10 @@ def sin( See Also -------- - :obj:`dpnp.arcsin` : Inverse sine, element-wise. + :obj:`dpnp.arcsin` : Trigonometric inverse sine, element-wise. + :obj:`dpnp.cos` : Trigonometric cosine, element-wise. + :obj:`dpnp.tan` : Trigonometric tangent, element-wise. :obj:`dpnp.sinh` : Hyperbolic sine, element-wise. - :obj:`dpnp.cos` : Cosine element-wise. Examples -------- @@ -1021,34 +1231,62 @@ def sin( ) -def sinh(x1): +def sinh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ - Trigonometric hyperbolic sine, element-wise. + Hyperbolic sine, element-wise. For full documentation refer to :obj:`numpy.sinh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic sine of each element of `x`. + Limitations ----------- - Input array is supported as :obj:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arcsinh` : Hyperbolic inverse sine, element-wise. + :obj:`dpnp.cosh` : Hyperbolic cosine, element-wise. + :obj:`dpnp.tanh` : Hyperbolic tangent, element-wise. + :obj:`dpnp.sin` : Trigonometric sine, element-wise. + Examples -------- >>> import dpnp as np >>> x = np.array([0, np.pi/2, np.pi]) - >>> out = np.sinh(x) - >>> [i for i in out] - [0.0, 2.3012989, 11.548739] + >>> np.sinh(x) + array([0.0, 2.3012989, 11.548739]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.sinh, + dpnp_sinh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_sinh(x1_desc).get_pyobj() - - return call_origin(numpy.sinh, x1, **kwargs) def sqrt( @@ -1069,7 +1307,7 @@ def sqrt( Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray An array of the same shape as `x`, containing the positive square-root of each element in `x`. If any element in `x` is complex, a complex array is returned (and the square-roots of @@ -1129,7 +1367,7 @@ def square( Returns ------- - y : dpnp.ndarray + out : dpnp.ndarray Element-wise `x * x`, of the same shape and dtype as `x`. Limitations @@ -1172,70 +1410,120 @@ def square( ) -def tan(x1, out=None, **kwargs): +def tan( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ - Compute tangent element-wise. + Trigonometric tangent, element-wise. For full documentation refer to :obj:`numpy.tan`. + Returns + ------- + out : dpnp.ndarray + The tangent of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. + :obj:`dpnp.sin` : Trigonometric sine, element-wise. + :obj:`dpnp.cos` : Trigonometric cosine, element-wise. + :obj:`dpnp.tanh` : Hyperbolic tangent, element-wise. + Examples -------- >>> import dpnp as np >>> x = np.array([-np.pi, np.pi/2, np.pi]) - >>> out = np.tan(x) - >>> [i for i in out] - [1.22460635e-16, 1.63317787e+16, -1.22460635e-16] + >>> np.tan(x) + array([1.22460635e-16, 1.63317787e+16, -1.22460635e-16]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.tan, + dpnp_tan, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_tan(x1_desc, out_desc).get_pyobj() - return call_origin(numpy.tan, x1, out=out, **kwargs) - -def tanh(x1): +def tanh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Compute hyperbolic tangent element-wise. For full documentation refer to :obj:`numpy.tanh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic tangent of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arctanh` : Hyperbolic inverse tangent, element-wise. + :obj:`dpnp.sinh` : Hyperbolic sine, element-wise. + :obj:`dpnp.cosh` : Hyperbolic cosine, element-wise. + :obj:`dpnp.tan` : Trigonometric tangent, element-wise. + Examples -------- >>> import dpnp as np - >>> x = np.array([-np.pi, np.pi/2, np.pi]) - >>> out = np.tanh(x) - >>> [i for i in out] - [-0.996272, 0.917152, 0.996272] + >>> x = np.array([0, -np.pi, np.pi/2, np.pi]) + >>> np.tanh(x) + array([0.0, -0.996272, 0.917152, 0.996272]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.tanh, + dpnp_tanh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_tanh(x1_desc).get_pyobj() - - return call_origin(numpy.tanh, x1, **kwargs) def unwrap(x1): diff --git a/tests/helper.py b/tests/helper.py index 42ada2949833..bb8ccd378a2b 100644 --- a/tests/helper.py +++ b/tests/helper.py @@ -7,7 +7,7 @@ import dpnp -def assert_dtype_allclose(dpnp_arr, numpy_arr): +def assert_dtype_allclose(dpnp_arr, numpy_arr, check_type=True): """ Assert DPNP and NumPy array based on maximum dtype resolution of input arrays for floating and complex types. @@ -24,7 +24,8 @@ def assert_dtype_allclose(dpnp_arr, numpy_arr): assert_allclose(dpnp_arr.asnumpy(), numpy_arr, atol=tol, rtol=tol) else: assert_array_equal(dpnp_arr.asnumpy(), numpy_arr) - assert dpnp_arr.dtype == numpy_arr.dtype + if check_type: + assert dpnp_arr.dtype == numpy_arr.dtype def get_complex_dtypes(device=None): diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index a78de781e3b0..128b7f7f5da2 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -34,7 +34,6 @@ tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_23_{axes=None, norm tests/third_party/intel/test_zero_copy_test1.py::test_dpnp_interaction_with_dpctl_memory -tests/test_arraymanipulation.py::TestHstack::test_generator tests/test_arraymanipulation.py::TestVstack::test_generator tests/test_linalg.py::test_cond[-1-[[1, 2, 3], [4, 5, 6], [7, 8, 9]]] @@ -84,6 +83,19 @@ tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: (dpnp tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp.asarray([(i, i) for i in x], [("a", object), ("b", dpnp.int32)])]] tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp.asarray(x).astype(dpnp.int8)] +tests/test_strides.py::test_strides_1arg[(10,)-None-cbrt] +tests/test_strides.py::test_strides_1arg[(10,)-None-degrees] +tests/test_strides.py::test_strides_1arg[(10,)-None-exp] +tests/test_strides.py::test_strides_1arg[(10,)-None-exp2] +tests/test_strides.py::test_strides_1arg[(10,)-None-expm1] +tests/test_strides.py::test_strides_1arg[(10,)-None-fabs] +tests/test_strides.py::test_strides_1arg[(10,)-None-log1p] +tests/test_strides.py::test_strides_1arg[(10,)-None-radians] +tests/test_strides.py::test_strides_1arg[(10,)-None-log2] +tests/test_strides.py::test_strides_1arg[(10,)-None-log10] +tests/test_strides.py::test_strides_erf[(10,)-None] +tests/test_strides.py::test_strides_reciprocal[(10,)-None] + tests/test_umath.py::test_umaths[('divmod', 'ii')] tests/test_umath.py::test_umaths[('divmod', 'll')] tests/test_umath.py::test_umaths[('divmod', 'ff')] @@ -111,10 +123,6 @@ tests/test_umath.py::test_umaths[('positive', 'i')] tests/test_umath.py::test_umaths[('positive', 'l')] tests/test_umath.py::test_umaths[('positive', 'f')] tests/test_umath.py::test_umaths[('positive', 'd')] -tests/test_umath.py::test_umaths[('rint', 'f')] -tests/test_umath.py::test_umaths[('rint', 'd')] -tests/test_umath.py::test_umaths[('signbit', 'f')] -tests/test_umath.py::test_umaths[('signbit', 'd')] tests/test_umath.py::test_umaths[('spacing', 'f')] tests/test_umath.py::test_umaths[('spacing', 'd')] @@ -161,14 +169,6 @@ tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestArrayCopyAn tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestArrayCopyAndView::test_view_non_contiguous_raise tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestNumPyArrayCopyView_param_0_{src_order='C'}::test_isinstance_numpy_view_copy_f tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestNumPyArrayCopyView_param_1_{src_order='F'}::test_isinstance_numpy_view_copy_f -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_0_{decimals=-3}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_1_{decimals=-2}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_2_{decimals=-1}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_3_{decimals=0}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_0_{decimals=-3}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_1_{decimals=-2}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_2_{decimals=-1}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_3_{decimals=0}::test_round_halfway_uint tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestArrayReduction::test_min_nan @@ -201,61 +201,17 @@ tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestCubReduction_pa tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestCubReduction_param_7_{order='F', shape=(10, 20, 30, 40)}::test_cub_max tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestCubReduction_param_7_{order='F', shape=(10, 20, 30, 40)}::test_cub_min tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_K_strides_reshape -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_contiguity2 tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_contiguity2_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_contiguity3 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_contiguity3_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_contiguity_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_full_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_ones_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_zeros_like_reshape_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_K_strides_reshape -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_contiguity2_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_contiguity3 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_contiguity3_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_contiguity_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_full_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_ones_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_zeros_like_reshape_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_K_strides_reshape -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_contiguity2_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_contiguity3 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_contiguity3_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_contiguity_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_full_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_ones_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_zeros_like_reshape_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_K_strides_reshape -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_contiguity2_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_contiguity3 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_contiguity3_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_contiguity_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_full_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_ones_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_zeros_like_reshape_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_K_strides_reshape -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_contiguity2_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_contiguity3 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_contiguity3_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_contiguity_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_full_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_ones_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_zeros_like_reshape_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_huge_size tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_huge_size_fill0 tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_int_huge_size tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_int_huge_size_fill0 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_like_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_like_contiguity3 tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_like_invalid_order tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_like_K_strides tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_like_subok @@ -265,8 +221,6 @@ tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_ones_like_s tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_zeros_like_subok tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_zeros_strides -tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_copy_order - tests/third_party/cupy/creation_tests/test_matrix.py::TestMatrix::test_diag_construction_from_list tests/third_party/cupy/creation_tests/test_matrix.py::TestMatrix::test_diag_construction_from_tuple tests/third_party/cupy/creation_tests/test_matrix.py::TestMatrix::test_diag_extraction_from_nested_list @@ -496,11 +450,7 @@ tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transpose tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transposed_tensordot_with_int_axes tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transposed_tensordot_with_list_axes tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_reversed_vdot -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_array_scalar -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_finite -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_infinite -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_infinite_equal_nan -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_min_int + tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_broadcast_not_allowed tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_diff_dtypes_is_equal tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_diff_dtypes_not_equal @@ -571,12 +521,10 @@ tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_frexp tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_ldexp tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_nextafter_combination tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_nextafter_float -tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_signbit tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_clip_min_max_none tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_external_clip4 tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_absolute_negative -tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_sign_negative tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_maximum_nan tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_minimum_nan tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_fmax_nan @@ -720,43 +668,7 @@ tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff_types[same] tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff_types[full] -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_0_{value=(14, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_0_{value=(14, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_1_{value=(15, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_1_{value=(15, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_2_{value=(16, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_2_{value=(16, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_3_{value=(14.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_3_{value=(14.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_4_{value=(15.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_4_{value=(15.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_5_{value=(16.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_5_{value=(16.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_0_{decimals=-100}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_0_{decimals=-100}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_1_{decimals=-99}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_1_{decimals=-99}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_2_{decimals=-90}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_2_{decimals=-90}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_3_{decimals=0}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_3_{decimals=0}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_4_{decimals=90}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_4_{decimals=90}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{decimals=99}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{decimals=99}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_6_{decimals=100}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_6_{decimals=100}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_around tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_fix -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_rint -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_rint_negative -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_round_ tests/third_party/cupy/math_tests/test_sumprod.py::TestSumprod::test_sum_out tests/third_party/cupy/math_tests/test_sumprod.py::TestSumprod::test_sum_out_wrong_shape tests/third_party/cupy/math_tests/test_sumprod.py::TestCumprod::test_ndarray_cumprod_2dim_with_axis diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl index 0b95dde9bb56..081f0b72350b 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -63,10 +63,6 @@ tests/test_umath.py::test_umaths[('positive', 'i')] tests/test_umath.py::test_umaths[('positive', 'l')] tests/test_umath.py::test_umaths[('positive', 'f')] tests/test_umath.py::test_umaths[('positive', 'd')] -tests/test_umath.py::test_umaths[('rint', 'f')] -tests/test_umath.py::test_umaths[('rint', 'd')] -tests/test_umath.py::test_umaths[('signbit', 'f')] -tests/test_umath.py::test_umaths[('signbit', 'd')] tests/test_umath.py::test_umaths[('spacing', 'f')] tests/test_umath.py::test_umaths[('spacing', 'd')] @@ -156,10 +152,6 @@ tests/third_party/cupy/core_tests/test_ndarray_conversion.py::TestNdarrayToBytes tests/third_party/cupy/core_tests/test_ndarray_conversion.py::TestNdarrayToBytes_param_2_{shape=(2, 3)}::test_item tests/third_party/cupy/core_tests/test_ndarray_conversion.py::TestNdarrayToBytes_param_3_{order='C', shape=(2, 3)}::test_item tests/third_party/cupy/core_tests/test_ndarray_conversion.py::TestNdarrayToBytes_param_4_{order='F', shape=(2, 3)}::test_item -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_0_{decimals=-3}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_1_{decimals=-2}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_2_{decimals=-1}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_3_{decimals=0}::test_round_halfway_float tests/third_party/cupy/creation_tests/test_matrix.py::TestMatrix::test_diag_construction tests/third_party/cupy/creation_tests/test_matrix.py::TestMatrix::test_diag_construction_from_list @@ -224,7 +216,6 @@ tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsMult tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsMultivariateNormal_param_3_{d=4, shape=(3, 2)}::test_normal tests/third_party/intel/test_zero_copy_test1.py::test_dpnp_interaction_with_dpctl_memory -tests/test_arraymanipulation.py::TestHstack::test_generator tests/test_arraymanipulation.py::TestVstack::test_generator tests/test_linalg.py::test_cond[-1-[[1, 2, 3], [4, 5, 6], [7, 8, 9]]] @@ -295,10 +286,6 @@ tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestArrayCopyAn tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestArrayCopyAndView::test_view_non_contiguous_raise tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestNumPyArrayCopyView_param_0_{src_order='C'}::test_isinstance_numpy_view_copy_f tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestNumPyArrayCopyView_param_1_{src_order='F'}::test_isinstance_numpy_view_copy_f -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_0_{decimals=-3}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_1_{decimals=-2}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_2_{decimals=-1}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_3_{decimals=0}::test_round_halfway_uint tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestArrayReduction::test_ptp_all tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestArrayReduction::test_ptp_all_keepdims @@ -329,61 +316,14 @@ tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestCubReduction_pa tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestCubReduction_param_7_{order='F', shape=(10, 20, 30, 40)}::test_cub_max tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestCubReduction_param_7_{order='F', shape=(10, 20, 30, 40)}::test_cub_min tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_K_strides_reshape -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_contiguity2_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_contiguity3 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_contiguity3_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_contiguity_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_empty_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_full_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_ones_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_0_{shape=4}::test_zeros_like_reshape_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_K_strides_reshape -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_contiguity2_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_contiguity3 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_contiguity3_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_contiguity_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_empty_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_full_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_ones_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_1_{shape=(4,)}::test_zeros_like_reshape_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_K_strides_reshape -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_contiguity2_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_contiguity3 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_contiguity3_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_contiguity_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_empty_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_full_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_ones_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_2_{shape=(4, 2)}::test_zeros_like_reshape_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_K_strides_reshape -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_contiguity2_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_contiguity3 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_contiguity3_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_contiguity_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_empty_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_full_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_ones_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_3_{shape=(4, 2, 3)}::test_zeros_like_reshape_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_K_strides_reshape -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_contiguity2_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_contiguity3 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_contiguity3_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_contiguity_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_empty_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_full_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_ones_like_reshape_cupy_only -tests/third_party/cupy/creation_tests/test_basic.py::TestBasicReshape_param_4_{shape=(5, 4, 2, 3)}::test_zeros_like_reshape_cupy_only tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_huge_size tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_huge_size_fill0 tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_int_huge_size tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_int_huge_size_fill0 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_like_contiguity2 -tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_like_contiguity3 tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_like_invalid_order tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_like_K_strides tests/third_party/cupy/creation_tests/test_basic.py::TestBasic::test_empty_like_subok @@ -430,8 +370,6 @@ tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_ tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_zero_num_no_endopoint_with_retstep tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_logspace_zero_num -tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_copy_order - tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_1_{axes=None, norm=None, s=(1, None), shape=(3, 4)}::test_fft2 tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_7_{axes=(), norm=None, s=None, shape=(3, 4)}::test_fft2 tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_7_{axes=(), norm=None, s=None, shape=(3, 4)}::test_ifft2 @@ -645,11 +583,7 @@ tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transpose tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_tensordot_zero_dim tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transposed_dot_with_out_f_contiguous tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transposed_multidim_vdot -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_array_scalar -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_finite -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_infinite -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_infinite_equal_nan -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_min_int + tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_broadcast_not_allowed tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_diff_dtypes_is_equal tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_diff_dtypes_not_equal @@ -716,7 +650,6 @@ tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_frexp tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_ldexp tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_nextafter_combination tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_nextafter_float -tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_signbit tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_clip_min_max_none tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_external_clip4 @@ -864,43 +797,7 @@ tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff_types[same] tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff_types[full] -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_0_{value=(14, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_0_{value=(14, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_1_{value=(15, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_1_{value=(15, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_2_{value=(16, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_2_{value=(16, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_3_{value=(14.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_3_{value=(14.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_4_{value=(15.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_4_{value=(15.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_5_{value=(16.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_5_{value=(16.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_0_{decimals=-100}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_0_{decimals=-100}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_1_{decimals=-99}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_1_{decimals=-99}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_2_{decimals=-90}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_2_{decimals=-90}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_3_{decimals=0}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_3_{decimals=0}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_4_{decimals=90}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_4_{decimals=90}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{decimals=99}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{decimals=99}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_6_{decimals=100}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_6_{decimals=100}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_around tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_fix -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_rint -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_rint_negative -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_round_ tests/third_party/cupy/math_tests/test_sumprod.py::TestCumprod::test_ndarray_cumprod_2dim_with_axis tests/third_party/cupy/math_tests/test_sumprod.py::TestDiff::test_diff_1dim tests/third_party/cupy/math_tests/test_sumprod.py::TestDiff::test_diff_1dim_with_n diff --git a/tests/skipped_tests_gpu_no_fp64.tbl b/tests/skipped_tests_gpu_no_fp64.tbl index 31c4b499cd0a..21aad4100b52 100644 --- a/tests/skipped_tests_gpu_no_fp64.tbl +++ b/tests/skipped_tests_gpu_no_fp64.tbl @@ -5,342 +5,6 @@ tests/test_linalg.py::test_norm1[None-3-[7]] tests/test_linalg.py::test_norm1[None-3-[1, 2]] tests/test_linalg.py::test_norm1[None-3-[1, 0]] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-bool_-bool_] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-bool_-int32] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-bool_-int64] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-bool_-None] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-int32-bool_] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-int32-int32] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-int32-int64] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-int32-None] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-int64-bool_] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-int64-int32] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-int64-int64] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-int64-None] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-None-bool_] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-None-int32] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-None-int64] -tests/test_mathematical.py::test_op_multiple_dtypes[[[1, 2], [3, 4]]-power-None-None] -tests/test_mathematical.py::TestMathematical::test_power[bool_-lhs0-rhs0] -tests/test_mathematical.py::TestMathematical::test_power[bool_-lhs0-rhs1] -tests/test_mathematical.py::TestMathematical::test_power[bool_-lhs0-3] -tests/test_mathematical.py::TestMathematical::test_power[bool_-lhs1-rhs0] -tests/test_mathematical.py::TestMathematical::test_power[bool_-lhs1-rhs1] -tests/test_mathematical.py::TestMathematical::test_power[bool_-lhs1-3] -tests/test_mathematical.py::TestMathematical::test_power[int32-lhs0-rhs0] -tests/test_mathematical.py::TestMathematical::test_power[int32-lhs0-rhs1] -tests/test_mathematical.py::TestMathematical::test_power[int32-lhs0-3] -tests/test_mathematical.py::TestMathematical::test_power[int32-lhs1-rhs0] -tests/test_mathematical.py::TestMathematical::test_power[int32-lhs1-rhs1] -tests/test_mathematical.py::TestMathematical::test_power[int32-lhs1-3] -tests/test_mathematical.py::TestMathematical::test_power[int64-lhs0-rhs0] -tests/test_mathematical.py::TestMathematical::test_power[int64-lhs0-rhs1] -tests/test_mathematical.py::TestMathematical::test_power[int64-lhs0-3] -tests/test_mathematical.py::TestMathematical::test_power[int64-lhs1-rhs0] -tests/test_mathematical.py::TestMathematical::test_power[int64-lhs1-rhs1] -tests/test_mathematical.py::TestMathematical::test_power[int64-lhs1-3] -tests/test_mathematical.py::TestMathematical::test_power[None-lhs0-rhs0] -tests/test_mathematical.py::TestMathematical::test_power[None-lhs0-3] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-0-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-0-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-0-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-0-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-0-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-0-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-0-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-0-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-1-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-1-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-1-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-1-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-1-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-1-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-1-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-1-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-5-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-5-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-5-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-5-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-5-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-5-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-5-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[0, 0], [0, 0]]-5-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-0-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-0-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-0-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-0-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-0-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-0-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-0-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-0-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-1-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-1-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-1-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-1-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-1-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-1-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-1-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-1-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-5-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-5-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-5-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-5-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-5-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-5-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-5-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [1, 2]]-5-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-0-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-0-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-0-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-0-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-0-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-0-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-0-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-0-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-1-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-1-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-1-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-1-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-1-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-1-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-1-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-1-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-5-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-5-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-5-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-5-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-5-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-5-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-5-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[1, 2], [3, 4]]-5-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-0-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-0-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-0-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-0-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-0-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-0-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-0-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-0-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-0-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-0-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-0-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-0-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-0-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-0-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-0-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-0-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-power-None-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-power-bool_-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-power-bool_-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-power-int32-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-power-int32-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-power-int64-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-power-int64-int] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-power-None-bool] -tests/test_mathematical.py::test_op_with_scalar[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-power-None-int] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-bool_-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-bool_-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-bool_-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-int32-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-int32-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-int32-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-int64-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-int64-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-int64-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-None-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-None-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1.5-None-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-bool_-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-bool_-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-bool_-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-int32-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-int32-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-int32-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-int64-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-int64-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-int64-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-None-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-None-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-1-None-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-bool_-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-bool_-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-bool_-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-int32-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-int32-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-int32-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-int64-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-int64-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-int64-int64] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-None-bool_] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-None-int32] -tests/test_mathematical.py::test_power[[[0, 0], [0, 0]]-5-None-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-bool_-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-bool_-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-bool_-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-int32-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-int32-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-int32-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-int64-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-int64-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-int64-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-None-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-None-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1.5-None-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-bool_-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-bool_-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-bool_-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-int32-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-int32-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-int32-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-int64-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-int64-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-int64-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-None-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-None-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-1-None-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-bool_-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-bool_-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-bool_-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-int32-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-int32-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-int32-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-int64-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-int64-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-int64-int64] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-None-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-None-int32] -tests/test_mathematical.py::test_power[[[1, 2], [1, 2]]-5-None-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-bool_-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-bool_-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-bool_-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-int32-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-int32-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-int32-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-int64-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-int64-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-int64-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-None-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-None-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1.5-None-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-bool_-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-bool_-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-bool_-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-int32-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-int32-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-int32-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-int64-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-int64-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-int64-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-None-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-None-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-1-None-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-bool_-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-bool_-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-bool_-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-int32-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-int32-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-int32-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-int64-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-int64-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-int64-int64] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-None-bool_] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-None-int32] -tests/test_mathematical.py::test_power[[[1, 2], [3, 4]]-5-None-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-bool_-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-bool_-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-bool_-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-int32-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-int32-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-int32-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-int64-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-int64-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-int64-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-None-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-None-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1.5-None-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-bool_-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-bool_-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-bool_-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-int32-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-int32-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-int32-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-int64-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-int64-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-int64-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-None-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-None-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-1-None-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-bool_-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-bool_-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-bool_-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-int32-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-int32-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-int32-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-int64-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-int64-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-int64-int64] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-None-bool_] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-None-int32] -tests/test_mathematical.py::test_power[[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]-5-None-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-bool_-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-bool_-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-bool_-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-int32-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-int32-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-int32-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-int64-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-int64-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-int64-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-None-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-None-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1.5-None-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-bool_-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-bool_-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-bool_-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-int32-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-int32-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-int32-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-int64-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-int64-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-int64-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-None-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-None-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-1-None-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-bool_-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-bool_-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-bool_-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-int32-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-int32-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-int32-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-int64-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-int64-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-int64-int64] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-None-bool_] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-None-int32] -tests/test_mathematical.py::test_power[[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]-5-None-int64] tests/test_mathematical.py::TestGradient::test_gradient_y1[array0] tests/test_mathematical.py::TestGradient::test_gradient_y1[array1] tests/test_mathematical.py::TestGradient::test_gradient_y1[array2] @@ -350,33 +14,8 @@ tests/test_mathematical.py::TestGradient::test_gradient_y1_dx[2-array2] tests/test_mathematical.py::TestGradient::test_gradient_y1_dx[3.5-array0] tests/test_mathematical.py::TestGradient::test_gradient_y1_dx[3.5-array1] tests/test_mathematical.py::TestGradient::test_gradient_y1_dx[3.5-array2] -tests/test_mathematical.py::TestPower::test_power[complex64] -tests/test_mathematical.py::TestPower::test_out_dtypes[bool_] -tests/test_mathematical.py::TestPower::test_out_dtypes[int32] -tests/test_mathematical.py::TestPower::test_out_dtypes[int64] -tests/test_mathematical.py::TestPower::test_out_dtypes[float32] -tests/test_mathematical.py::TestPower::test_out_overlap[int32] -tests/test_mathematical.py::TestPower::test_out_overlap[int64] -tests/test_mathematical.py::TestPower::test_inplace_strided_out[int32] -tests/test_mathematical.py::TestPower::test_inplace_strided_out[int64] -tests/test_mathematical.py::TestPower::test_invalid_shape[(0,)] -tests/test_mathematical.py::TestPower::test_invalid_shape[(15, )] -tests/test_mathematical.py::TestPower::test_invalid_shape[(2,2)] -tests/test_mathematical.py::TestPower::test_integer_power_of_0_or_1[int32-0] -tests/test_mathematical.py::TestPower::test_integer_power_of_0_or_1[int32-1] -tests/test_mathematical.py::TestPower::test_integer_power_of_0_or_1[int64-0] -tests/test_mathematical.py::TestPower::test_integer_power_of_0_or_1[int64-1] -tests/test_mathematical.py::TestPower::test_integer_to_negative_power[int32] -tests/test_mathematical.py::TestPower::test_integer_to_negative_power[int64] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arccos] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arccosh] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arcsin] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arcsinh] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arctan] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arctanh] tests/test_strides.py::test_strides_1arg[(10,)-int32-cbrt] -tests/test_strides.py::test_strides_1arg[(10,)-int32-cosh] tests/test_strides.py::test_strides_1arg[(10,)-int32-degrees] tests/test_strides.py::test_strides_1arg[(10,)-int32-exp] tests/test_strides.py::test_strides_1arg[(10,)-int32-exp2] @@ -386,16 +25,7 @@ tests/test_strides.py::test_strides_1arg[(10,)-int32-log10] tests/test_strides.py::test_strides_1arg[(10,)-int32-log1p] tests/test_strides.py::test_strides_1arg[(10,)-int32-log2] tests/test_strides.py::test_strides_1arg[(10,)-int32-radians] -tests/test_strides.py::test_strides_1arg[(10,)-int32-sinh] -tests/test_strides.py::test_strides_1arg[(10,)-int32-tanh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arccos] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arccosh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arcsin] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arcsinh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arctan] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arctanh] tests/test_strides.py::test_strides_1arg[(10,)-int64-cbrt] -tests/test_strides.py::test_strides_1arg[(10,)-int64-cosh] tests/test_strides.py::test_strides_1arg[(10,)-int64-degrees] tests/test_strides.py::test_strides_1arg[(10,)-int64-exp] tests/test_strides.py::test_strides_1arg[(10,)-int64-exp2] @@ -405,16 +35,7 @@ tests/test_strides.py::test_strides_1arg[(10,)-int64-log10] tests/test_strides.py::test_strides_1arg[(10,)-int64-log1p] tests/test_strides.py::test_strides_1arg[(10,)-int64-log2] tests/test_strides.py::test_strides_1arg[(10,)-int64-radians] -tests/test_strides.py::test_strides_1arg[(10,)-int64-sinh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-tanh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arccos] -tests/test_strides.py::test_strides_1arg[(10,)-None-arccosh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arcsin] -tests/test_strides.py::test_strides_1arg[(10,)-None-arcsinh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arctan] -tests/test_strides.py::test_strides_1arg[(10,)-None-arctanh] tests/test_strides.py::test_strides_1arg[(10,)-None-cbrt] -tests/test_strides.py::test_strides_1arg[(10,)-None-cosh] tests/test_strides.py::test_strides_1arg[(10,)-None-degrees] tests/test_strides.py::test_strides_1arg[(10,)-None-exp] tests/test_strides.py::test_strides_1arg[(10,)-None-exp2] @@ -424,43 +45,8 @@ tests/test_strides.py::test_strides_1arg[(10,)-None-log10] tests/test_strides.py::test_strides_1arg[(10,)-None-log1p] tests/test_strides.py::test_strides_1arg[(10,)-None-log2] tests/test_strides.py::test_strides_1arg[(10,)-None-radians] -tests/test_strides.py::test_strides_1arg[(10,)-None-sinh] -tests/test_strides.py::test_strides_1arg[(10,)-None-tanh] -tests/test_strides.py::test_strides_tan[(10,)-int32] -tests/test_strides.py::test_strides_tan[(10,)-int64] -tests/test_strides.py::test_strides_tan[(10,)-None] -tests/test_strides.py::test_strides_2args[(3, 3)-int32-power] -tests/test_strides.py::test_strides_2args[(3, 3)-int64-power] -tests/test_strides.py::test_strides_2args[(3, 3)-None-power] -tests/test_strides.py::test_strided_out_2args[int32-power] -tests/test_strides.py::test_strided_out_2args[int64-power] -tests/test_strides.py::test_strided_out_2args[float32-power] -tests/test_strides.py::test_strided_out_2args[None-power] -tests/test_strides.py::test_strided_in_out_2args[int32-power] -tests/test_strides.py::test_strided_in_out_2args[int64-power] - -tests/test_sycl_queue.py::test_array_creation[opencl:gpu:0-arange-arg0-kwargs0] -tests/test_sycl_queue.py::test_array_creation[level_zero:gpu:0-arange-arg0-kwargs0] -tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-gradient-data10] -tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-gradient-data10] -tests/test_sycl_queue.py::test_2in_1out[opencl:gpu:0-power-data112-data212] -tests/test_sycl_queue.py::test_2in_1out[level_zero:gpu:0-power-data112-data212] -tests/test_sycl_queue.py::test_out_2in_1out[opencl:gpu:0-power-data19-data29] -tests/test_sycl_queue.py::test_out_2in_1out[level_zero:gpu:0-power-data19-data29] -tests/test_sycl_queue.py::test_eig[opencl:gpu:0] -tests/test_sycl_queue.py::test_eig[level_zero:gpu:0] -tests/test_sycl_queue.py::test_eigh[opencl:gpu:0] -tests/test_sycl_queue.py::test_eigh[level_zero:gpu:0] -tests/test_sycl_queue.py::test_qr[opencl:gpu:0] -tests/test_sycl_queue.py::test_qr[level_zero:gpu:0] -tests/test_sycl_queue.py::test_svd[opencl:gpu:0] -tests/test_sycl_queue.py::test_svd[level_zero:gpu:0] -tests/test_sycl_queue.py::test_to_device[opencl:gpu:0-opencl:cpu:0] -tests/test_sycl_queue.py::test_to_device[level_zero:gpu:0-opencl:cpu:0] tests/test_umath.py::test_umaths[('floor_divide', 'ff')] -tests/test_umath.py::test_umaths[('power', 'ii')] -tests/test_umath.py::test_umaths[('power', 'll')] tests/test_umath.py::TestExp::test_exp tests/test_umath.py::TestExp::test_invalid_dtype[numpy.float32] @@ -469,323 +55,12 @@ tests/test_umath.py::TestExp::test_invalid_dtype[numpy.int32] tests/test_umath.py::TestExp::test_invalid_shape[(0,)] tests/test_umath.py::TestExp::test_invalid_shape[(15, )] tests/test_umath.py::TestExp::test_invalid_shape[(2,2)] -tests/test_umath.py::TestArcsin::test_arcsin -tests/test_umath.py::TestArcsin::test_invalid_dtype[numpy.float32] -tests/test_umath.py::TestArcsin::test_invalid_dtype[numpy.int64] -tests/test_umath.py::TestArcsin::test_invalid_dtype[numpy.int32] -tests/test_umath.py::TestArcsin::test_invalid_shape[(0,)] -tests/test_umath.py::TestArcsin::test_invalid_shape[(15, )] -tests/test_umath.py::TestArcsin::test_invalid_shape[(2,2)] -tests/test_umath.py::TestArctan::test_arctan -tests/test_umath.py::TestArctan::test_invalid_dtype[numpy.float32] -tests/test_umath.py::TestArctan::test_invalid_dtype[numpy.int64] -tests/test_umath.py::TestArctan::test_invalid_dtype[numpy.int32] -tests/test_umath.py::TestArctan::test_invalid_shape[(0,)] -tests/test_umath.py::TestArctan::test_invalid_shape[(15, )] -tests/test_umath.py::TestArctan::test_invalid_shape[(2,2)] -tests/test_umath.py::TestTan::test_tan -tests/test_umath.py::TestTan::test_invalid_dtype[numpy.float32] -tests/test_umath.py::TestTan::test_invalid_dtype[numpy.int64] -tests/test_umath.py::TestTan::test_invalid_dtype[numpy.int32] -tests/test_umath.py::TestTan::test_invalid_shape[(0,)] -tests/test_umath.py::TestTan::test_invalid_shape[(15, )] -tests/test_umath.py::TestTan::test_invalid_shape[(2,2)] -tests/test_umath.py::TestArctan2::test_arctan2 -tests/test_umath.py::TestArctan2::test_invalid_shape[(0,)] -tests/test_umath.py::TestArctan2::test_invalid_shape[(15, )] -tests/test_umath.py::TestArctan2::test_invalid_shape[(2,2)] tests/test_umath.py::TestSqrt::test_sqrt_complex[complex64] -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRound_param_2_{decimals=0}::test_round -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRound_param_0_{decimals=-2}::test_round_out -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRound_param_1_{decimals=-1}::test_round_out -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRound_param_2_{decimals=0}::test_round_out -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRound_param_3_{decimals=1}::test_round_out -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRound_param_4_{decimals=2}::test_round_out - -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_arange_no_dtype_float -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_array_start_stop -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_float_overflow -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_float_underflow -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_mixed_start_stop -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_mixed_start_stop2 -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_no_dtype_float -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_no_dtype_int -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_linspace_start_stop_list -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_logspace_base -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_logspace_no_dtype_float -tests/third_party/cupy/creation_tests/test_ranges.py::TestRanges::test_logspace_no_dtype_int -tests/third_party/cupy/creation_tests/test_ranges.py::TestMgrid::test_mgrid2 -tests/third_party/cupy/creation_tests/test_ranges.py::TestMgrid::test_mgrid4 -tests/third_party/cupy/creation_tests/test_ranges.py::TestOgrid::test_ogrid2 - -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_2_{n=None, norm=None, shape=(10,)}::test_fft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_2_{n=None, norm=None, shape=(10,)}::test_ifft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_3_{n=None, norm=None, shape=(10, 10)}::test_fft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_3_{n=None, norm=None, shape=(10, 10)}::test_ifft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_6_{n=None, norm='ortho', shape=(10,)}::test_fft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_6_{n=None, norm='ortho', shape=(10,)}::test_ifft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_7_{n=None, norm='ortho', shape=(10, 10)}::test_fft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_7_{n=None, norm='ortho', shape=(10, 10)}::test_ifft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_0_{axes=None, norm=None, s=None, shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_0_{axes=None, norm=None, s=None, shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_2_{axes=None, norm=None, s=(1, 5), shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_2_{axes=None, norm=None, s=(1, 5), shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_3_{axes=(-2, -1), norm=None, s=None, shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_3_{axes=(-2, -1), norm=None, s=None, shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_4_{axes=(-1, -2), norm=None, s=None, shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_4_{axes=(-1, -2), norm=None, s=None, shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_5_{axes=(0,), norm=None, s=None, shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_5_{axes=(0,), norm=None, s=None, shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_6_{axes=None, norm='ortho', s=None, shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_6_{axes=None, norm='ortho', s=None, shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_8_{axes=None, norm=None, s=None, shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_8_{axes=None, norm=None, s=None, shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_10_{axes=None, norm=None, s=(1, 4, 10), shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_10_{axes=None, norm=None, s=(1, 4, 10), shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_11_{axes=(-3, -2, -1), norm=None, s=None, shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_11_{axes=(-3, -2, -1), norm=None, s=None, shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_12_{axes=(-1, -2, -3), norm=None, s=None, shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_12_{axes=(-1, -2, -3), norm=None, s=None, shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_13_{axes=(0, 1), norm=None, s=None, shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_13_{axes=(0, 1), norm=None, s=None, shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_14_{axes=None, norm='ortho', s=None, shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_14_{axes=None, norm='ortho', s=None, shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_17_{axes=None, norm=None, s=None, shape=(2, 3, 4, 5)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_17_{axes=None, norm=None, s=None, shape=(2, 3, 4, 5)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_0_{axes=None, norm=None, s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_0_{axes=None, norm=None, s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_2_{axes=None, norm=None, s=(1, 5), shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_2_{axes=None, norm=None, s=(1, 5), shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_3_{axes=(-2, -1), norm=None, s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_3_{axes=(-2, -1), norm=None, s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_4_{axes=(-1, -2), norm=None, s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_4_{axes=(-1, -2), norm=None, s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_5_{axes=[-1, -2], norm=None, s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_5_{axes=[-1, -2], norm=None, s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_6_{axes=(0,), norm=None, s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_6_{axes=(0,), norm=None, s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_8_{axes=None, norm='ortho', s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_8_{axes=None, norm='ortho', s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_9_{axes=None, norm=None, s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_9_{axes=None, norm=None, s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_11_{axes=None, norm=None, s=(1, 4, 10), shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_11_{axes=None, norm=None, s=(1, 4, 10), shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_12_{axes=(-3, -2, -1), norm=None, s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_12_{axes=(-3, -2, -1), norm=None, s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_13_{axes=(-1, -2, -3), norm=None, s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_13_{axes=(-1, -2, -3), norm=None, s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_14_{axes=(-1, -3), norm=None, s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_14_{axes=(-1, -3), norm=None, s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_15_{axes=(0, 1), norm=None, s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_15_{axes=(0, 1), norm=None, s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_16_{axes=None, norm='ortho', s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_16_{axes=None, norm='ortho', s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_19_{axes=(2, 0, 1), norm='ortho', s=(4, 3, 2), shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_19_{axes=(2, 0, 1), norm='ortho', s=(4, 3, 2), shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_20_{axes=None, norm=None, s=None, shape=(2, 3, 4, 5)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_20_{axes=None, norm=None, s=None, shape=(2, 3, 4, 5)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_0_{n=None, norm=None, shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_0_{n=None, norm=None, shape=(10,)}::test_rfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_1_{n=None, norm=None, shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_1_{n=None, norm=None, shape=(10, 10)}::test_rfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_2_{n=None, norm='ortho', shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_3_{n=None, norm='ortho', shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_4_{n=5, norm=None, shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_5_{n=5, norm=None, shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_6_{n=5, norm='ortho', shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_7_{n=5, norm='ortho', shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_8_{n=10, norm=None, shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_9_{n=10, norm=None, shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_10_{n=10, norm='ortho', shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_11_{n=10, norm='ortho', shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_12_{n=15, norm=None, shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_13_{n=15, norm=None, shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_14_{n=15, norm='ortho', shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_15_{n=15, norm='ortho', shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_0_{n=None, norm=None, shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_0_{n=None, norm=None, shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_1_{n=None, norm=None, shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_1_{n=None, norm=None, shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_2_{n=None, norm='ortho', shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_2_{n=None, norm='ortho', shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_3_{n=None, norm='ortho', shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_3_{n=None, norm='ortho', shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_4_{n=5, norm=None, shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_4_{n=5, norm=None, shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_5_{n=5, norm=None, shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_5_{n=5, norm=None, shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_6_{n=5, norm='ortho', shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_6_{n=5, norm='ortho', shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_7_{n=5, norm='ortho', shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_7_{n=5, norm='ortho', shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_8_{n=10, norm=None, shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_8_{n=10, norm=None, shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_9_{n=10, norm=None, shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_9_{n=10, norm=None, shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_10_{n=10, norm='ortho', shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_10_{n=10, norm='ortho', shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_11_{n=10, norm='ortho', shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_11_{n=10, norm='ortho', shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_12_{n=15, norm=None, shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_12_{n=15, norm=None, shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_13_{n=15, norm=None, shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_13_{n=15, norm=None, shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_14_{n=15, norm='ortho', shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_14_{n=15, norm='ortho', shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_15_{n=15, norm='ortho', shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_15_{n=15, norm='ortho', shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_0_{d=1, n=1}::test_fftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_0_{d=1, n=1}::test_rfftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_1_{d=0.5, n=10}::test_fftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_1_{d=0.5, n=10}::test_rfftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_2_{d=2, n=100}::test_fftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_2_{d=2, n=100}::test_rfftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_2_{n=None, norm=None, shape=(10,)}::test_fft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_2_{n=None, norm=None, shape=(10,)}::test_ifft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_3_{n=None, norm=None, shape=(10, 10)}::test_fft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_3_{n=None, norm=None, shape=(10, 10)}::test_ifft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_6_{n=None, norm='ortho', shape=(10,)}::test_fft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_6_{n=None, norm='ortho', shape=(10,)}::test_ifft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_7_{n=None, norm='ortho', shape=(10, 10)}::test_fft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft_param_7_{n=None, norm='ortho', shape=(10, 10)}::test_ifft -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_0_{axes=None, norm=None, s=None, shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_0_{axes=None, norm=None, s=None, shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_2_{axes=None, norm=None, s=(1, 5), shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_2_{axes=None, norm=None, s=(1, 5), shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_3_{axes=(-2, -1), norm=None, s=None, shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_3_{axes=(-2, -1), norm=None, s=None, shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_4_{axes=(-1, -2), norm=None, s=None, shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_4_{axes=(-1, -2), norm=None, s=None, shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_5_{axes=(0,), norm=None, s=None, shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_5_{axes=(0,), norm=None, s=None, shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_6_{axes=None, norm='ortho', s=None, shape=(3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_6_{axes=None, norm='ortho', s=None, shape=(3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_8_{axes=None, norm=None, s=None, shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_8_{axes=None, norm=None, s=None, shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_10_{axes=None, norm=None, s=(1, 4, 10), shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_10_{axes=None, norm=None, s=(1, 4, 10), shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_11_{axes=(-3, -2, -1), norm=None, s=None, shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_11_{axes=(-3, -2, -1), norm=None, s=None, shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_12_{axes=(-1, -2, -3), norm=None, s=None, shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_12_{axes=(-1, -2, -3), norm=None, s=None, shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_13_{axes=(0, 1), norm=None, s=None, shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_13_{axes=(0, 1), norm=None, s=None, shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_14_{axes=None, norm='ortho', s=None, shape=(2, 3, 4)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_14_{axes=None, norm='ortho', s=None, shape=(2, 3, 4)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_17_{axes=None, norm=None, s=None, shape=(2, 3, 4, 5)}::test_fft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_17_{axes=None, norm=None, s=None, shape=(2, 3, 4, 5)}::test_ifft2 -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_0_{axes=None, norm=None, s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_0_{axes=None, norm=None, s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_2_{axes=None, norm=None, s=(1, 5), shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_2_{axes=None, norm=None, s=(1, 5), shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_3_{axes=(-2, -1), norm=None, s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_3_{axes=(-2, -1), norm=None, s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_4_{axes=(-1, -2), norm=None, s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_4_{axes=(-1, -2), norm=None, s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_5_{axes=[-1, -2], norm=None, s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_5_{axes=[-1, -2], norm=None, s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_6_{axes=(0,), norm=None, s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_6_{axes=(0,), norm=None, s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_8_{axes=None, norm='ortho', s=None, shape=(3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_8_{axes=None, norm='ortho', s=None, shape=(3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_9_{axes=None, norm=None, s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_9_{axes=None, norm=None, s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_11_{axes=None, norm=None, s=(1, 4, 10), shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_11_{axes=None, norm=None, s=(1, 4, 10), shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_12_{axes=(-3, -2, -1), norm=None, s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_12_{axes=(-3, -2, -1), norm=None, s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_13_{axes=(-1, -2, -3), norm=None, s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_13_{axes=(-1, -2, -3), norm=None, s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_14_{axes=(-1, -3), norm=None, s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_14_{axes=(-1, -3), norm=None, s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_15_{axes=(0, 1), norm=None, s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_15_{axes=(0, 1), norm=None, s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_16_{axes=None, norm='ortho', s=None, shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_16_{axes=None, norm='ortho', s=None, shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_19_{axes=(2, 0, 1), norm='ortho', s=(4, 3, 2), shape=(2, 3, 4)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_19_{axes=(2, 0, 1), norm='ortho', s=(4, 3, 2), shape=(2, 3, 4)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_20_{axes=None, norm=None, s=None, shape=(2, 3, 4, 5)}::test_fftn -tests/third_party/cupy/fft_tests/test_fft.py::TestFftn_param_20_{axes=None, norm=None, s=None, shape=(2, 3, 4, 5)}::test_ifftn -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_0_{n=None, norm=None, shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_0_{n=None, norm=None, shape=(10,)}::test_rfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_1_{n=None, norm=None, shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_1_{n=None, norm=None, shape=(10, 10)}::test_rfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_2_{n=None, norm='ortho', shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_3_{n=None, norm='ortho', shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_4_{n=5, norm=None, shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_5_{n=5, norm=None, shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_6_{n=5, norm='ortho', shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_7_{n=5, norm='ortho', shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_8_{n=10, norm=None, shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_9_{n=10, norm=None, shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_10_{n=10, norm='ortho', shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_11_{n=10, norm='ortho', shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_12_{n=15, norm=None, shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_13_{n=15, norm=None, shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_14_{n=15, norm='ortho', shape=(10,)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestRfft_param_15_{n=15, norm='ortho', shape=(10, 10)}::test_irfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_0_{n=None, norm=None, shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_0_{n=None, norm=None, shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_1_{n=None, norm=None, shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_1_{n=None, norm=None, shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_2_{n=None, norm='ortho', shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_2_{n=None, norm='ortho', shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_3_{n=None, norm='ortho', shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_3_{n=None, norm='ortho', shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_4_{n=5, norm=None, shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_4_{n=5, norm=None, shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_5_{n=5, norm=None, shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_5_{n=5, norm=None, shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_6_{n=5, norm='ortho', shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_6_{n=5, norm='ortho', shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_7_{n=5, norm='ortho', shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_7_{n=5, norm='ortho', shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_8_{n=10, norm=None, shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_8_{n=10, norm=None, shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_9_{n=10, norm=None, shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_9_{n=10, norm=None, shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_10_{n=10, norm='ortho', shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_10_{n=10, norm='ortho', shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_11_{n=10, norm='ortho', shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_11_{n=10, norm='ortho', shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_12_{n=15, norm=None, shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_12_{n=15, norm=None, shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_13_{n=15, norm=None, shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_13_{n=15, norm=None, shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_14_{n=15, norm='ortho', shape=(10,)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_14_{n=15, norm='ortho', shape=(10,)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_15_{n=15, norm='ortho', shape=(10, 10)}::test_hfft -tests/third_party/cupy/fft_tests/test_fft.py::TestHfft_param_15_{n=15, norm='ortho', shape=(10, 10)}::test_ihfft -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_0_{d=1, n=1}::test_fftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_0_{d=1, n=1}::test_rfftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_1_{d=0.5, n=10}::test_fftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_1_{d=0.5, n=10}::test_rfftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_2_{d=2, n=100}::test_fftfreq -tests/third_party/cupy/fft_tests/test_fft.py::TestFftfreq_param_2_{d=2, n=100}::test_rfftfreq - -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalue_param_0_{UPLO='U'}::test_eigh tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalue_param_0_{UPLO='U'}::test_eigh_batched -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalue_param_1_{UPLO='L'}::test_eigh tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalue_param_1_{UPLO='L'}::test_eigh_batched -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_0_{UPLO='U', shape=(0, 0)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_1_{UPLO='U', shape=(2, 0, 0)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_2_{UPLO='U', shape=(0, 3, 3)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_3_{UPLO='L', shape=(0, 0)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_4_{UPLO='L', shape=(2, 0, 0)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_5_{UPLO='L', shape=(0, 3, 3)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalue_param_0_{UPLO='U'}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalue_param_0_{UPLO='U'}::test_eigh_batched -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalue_param_1_{UPLO='L'}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalue_param_1_{UPLO='L'}::test_eigh_batched -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_0_{UPLO='U', shape=(0, 0)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_1_{UPLO='U', shape=(2, 0, 0)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_2_{UPLO='U', shape=(0, 3, 3)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_3_{UPLO='L', shape=(0, 0)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_4_{UPLO='L', shape=(2, 0, 0)}::test_eigh -tests/third_party/cupy/linalg_tests/test_eigenvalue.py::TestEigenvalueEmpty_param_5_{UPLO='L', shape=(0, 3, 3)}::test_eigh -tests/third_party/cupy/linalg_tests/test_einsum.py::TestEinSumLarge_param_0_{opt=('greedy', 0), subscript='a,b,c->abc'}::test_einsum + tests/third_party/cupy/linalg_tests/test_product.py::TestDot_param_0_{shape=((2, 3, 4), (3, 4, 2)), trans_a=True, trans_b=True}::test_dot tests/third_party/cupy/linalg_tests/test_product.py::TestDot_param_1_{shape=((2, 3, 4), (3, 4, 2)), trans_a=True, trans_b=False}::test_dot tests/third_party/cupy/linalg_tests/test_product.py::TestDot_param_2_{shape=((2, 3, 4), (3, 4, 2)), trans_a=False, trans_b=True}::test_dot @@ -870,134 +145,9 @@ tests/third_party/cupy/linalg_tests/test_product.py::TestDotFor0Dim_param_8_{sha tests/third_party/cupy/linalg_tests/test_product.py::TestDotFor0Dim_param_9_{shape=((4, 2), ()), trans_a=True, trans_b=False}::test_dot tests/third_party/cupy/linalg_tests/test_product.py::TestDotFor0Dim_param_10_{shape=((4, 2), ()), trans_a=False, trans_b=True}::test_dot tests/third_party/cupy/linalg_tests/test_product.py::TestDotFor0Dim_param_11_{shape=((4, 2), ()), trans_a=False, trans_b=False}::test_dot -tests/third_party/cupy/linalg_tests/test_product.py::TestCrossProduct_param_0_{params=((3,), (3,), -1, -1, -1)}::test_cross -tests/third_party/cupy/linalg_tests/test_product.py::TestCrossProduct_param_1_{params=((1, 2), (1, 2), -1, -1, 1)}::test_cross -tests/third_party/cupy/linalg_tests/test_product.py::TestCrossProduct_param_2_{params=((1, 3), (1, 3), 1, -1, -1)}::test_cross -tests/third_party/cupy/linalg_tests/test_product.py::TestCrossProduct_param_3_{params=((1, 2), (1, 3), -1, -1, 1)}::test_cross -tests/third_party/cupy/linalg_tests/test_product.py::TestCrossProduct_param_4_{params=((2, 2), (1, 3), -1, -1, 0)}::test_cross -tests/third_party/cupy/linalg_tests/test_product.py::TestCrossProduct_param_5_{params=((3, 3), (1, 2), 0, -1, -1)}::test_cross -tests/third_party/cupy/linalg_tests/test_product.py::TestCrossProduct_param_6_{params=((0, 3), (0, 3), -1, -1, -1)}::test_cross -tests/third_party/cupy/linalg_tests/test_product.py::TestCrossProduct_param_7_{params=((2, 0, 3), (2, 0, 3), 0, 0, 0)}::test_cross -tests/third_party/cupy/linalg_tests/test_product.py::TestCrossProduct_param_8_{params=((2, 4, 5, 3), (2, 4, 5, 3), -1, -1, 0)}::test_cross -tests/third_party/cupy/linalg_tests/test_product.py::TestCrossProduct_param_9_{params=((2, 4, 5, 2), (2, 4, 5, 2), 0, 0, -1)}::test_cross - -tests/third_party/cupy/logic_tests/test_content.py::TestContent::test_isfinite -tests/third_party/cupy/logic_tests/test_content.py::TestContent::test_isinf -tests/third_party/cupy/logic_tests/test_content.py::TestContent::test_isnan tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticUnary_param_29_{arg1=2.0, name='reciprocal'}::test_unary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_6_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0., 1., 2.], [3., 4., 5.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_10_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_10_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_10_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_10_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_14_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_14_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_14_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_14_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_38_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0., 1., 2.], [3., 4., 5.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_42_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_42_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_42_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_42_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_46_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_46_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_46_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_46_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_66_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_66_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_66_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[6., 5., 4.], [3., 2., 1.]], dtype=float32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_66_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[6., 5., 4.], [3., 2., 1.]], dtype=float32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_70_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0., 1., 2.], [3., 4., 5.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_70_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_70_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[6., 5., 4.], [3., 2., 1.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_70_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[6., 5., 4.], [3., 2., 1.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_74_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_74_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_74_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_74_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_78_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_78_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_78_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_78_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_82_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=0, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_82_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=0, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_86_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=0.0, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_90_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=2, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_90_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=2, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_94_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=2.0, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_98_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_98_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_98_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[6., 5., 4.], [3., 2., 1.]], dtype=float32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_98_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[6., 5., 4.], [3., 2., 1.]], dtype=float32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_102_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_102_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0., 1., 2.], [3., 4., 5.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_102_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[6., 5., 4.], [3., 2., 1.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_102_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[6., 5., 4.], [3., 2., 1.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_106_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_106_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_106_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_106_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_110_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_110_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_110_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_110_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_114_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=0, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_114_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=0, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_118_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=0.0, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_122_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=2, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_122_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=2, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_126_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=2.0, name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_134_{arg1=0, arg2=array([[0., 1., 2.], [3., 4., 5.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_138_{arg1=0, arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_138_{arg1=0, arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_138_{arg1=0, arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_138_{arg1=0, arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_142_{arg1=0, arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_142_{arg1=0, arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_142_{arg1=0, arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_142_{arg1=0, arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_166_{arg1=0.0, arg2=array([[0., 1., 2.], [3., 4., 5.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_170_{arg1=0.0, arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_174_{arg1=0.0, arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_198_{arg1=2, arg2=array([[0., 1., 2.], [3., 4., 5.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_202_{arg1=2, arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_202_{arg1=2, arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_202_{arg1=2, arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_202_{arg1=2, arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_206_{arg1=2, arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_206_{arg1=2, arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_206_{arg1=2, arg2=array([[6, 5, 4], [3, 2, 1]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_206_{arg1=2, arg2=array([[6, 5, 4], [3, 2, 1]], dtype=int64), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_230_{arg1=2.0, arg2=array([[0., 1., 2.], [3., 4., 5.]]), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_234_{arg1=2.0, arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), name='power'}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary_param_238_{arg1=2.0, arg2=array([[0, 1, 2], [3, 4, 5]]), name='power'}::test_binary - -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_128_{arg1=array([-3., -2., -1., 1., 2., 3.], dtype=float32), arg2=array([-3., -2., -1., 1., 2., 3.], dtype=float32), dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_134_{arg1=array([-3., -2., -1., 1., 2., 3.], dtype=float32), arg2=array([-3., -2., -1., 1., 2., 3.]), dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_135_{arg1=array([-3., -2., -1., 1., 2., 3.], dtype=float32), arg2=array([-3., -2., -1., 1., 2., 3.]), dtype=float64, name='power', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_140_{arg1=array([-3., -2., -1., 1., 2., 3.], dtype=float32), arg2=0.0, dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_146_{arg1=array([-3., -2., -1., 1., 2., 3.], dtype=float32), arg2=2.0, dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_152_{arg1=array([-3., -2., -1., 1., 2., 3.], dtype=float32), arg2=-2.0, dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_158_{arg1=array([-3., -2., -1., 1., 2., 3.]), arg2=array([-3., -2., -1., 1., 2., 3.], dtype=float32), dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_159_{arg1=array([-3., -2., -1., 1., 2., 3.]), arg2=array([-3., -2., -1., 1., 2., 3.], dtype=float32), dtype=float64, name='power', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_164_{arg1=array([-3., -2., -1., 1., 2., 3.]), arg2=array([-3., -2., -1., 1., 2., 3.]), dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_165_{arg1=array([-3., -2., -1., 1., 2., 3.]), arg2=array([-3., -2., -1., 1., 2., 3.]), dtype=float64, name='power', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_170_{arg1=array([-3., -2., -1., 1., 2., 3.]), arg2=0.0, dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_171_{arg1=array([-3., -2., -1., 1., 2., 3.]), arg2=0.0, dtype=float64, name='power', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_176_{arg1=array([-3., -2., -1., 1., 2., 3.]), arg2=2.0, dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_177_{arg1=array([-3., -2., -1., 1., 2., 3.]), arg2=2.0, dtype=float64, name='power', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_183_{arg1=array([-3., -2., -1., 1., 2., 3.]), arg2=-2.0, dtype=float64, name='power', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_188_{arg1=0.0, arg2=array([-3., -2., -1., 1., 2., 3.], dtype=float32), dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_194_{arg1=0.0, arg2=array([-3., -2., -1., 1., 2., 3.]), dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_195_{arg1=0.0, arg2=array([-3., -2., -1., 1., 2., 3.]), dtype=float64, name='power', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_218_{arg1=2.0, arg2=array([-3., -2., -1., 1., 2., 3.], dtype=float32), dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_224_{arg1=2.0, arg2=array([-3., -2., -1., 1., 2., 3.]), dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_225_{arg1=2.0, arg2=array([-3., -2., -1., 1., 2., 3.]), dtype=float64, name='power', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_248_{arg1=-2.0, arg2=array([-3., -2., -1., 1., 2., 3.], dtype=float32), dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_254_{arg1=-2.0, arg2=array([-3., -2., -1., 1., 2., 3.]), dtype=float64, name='power', use_dtype=True}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_255_{arg1=-2.0, arg2=array([-3., -2., -1., 1., 2., 3.]), dtype=float64, name='power', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_280_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='fmod', use_dtype=True}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_288_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='fmod', use_dtype=True}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_289_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='fmod', use_dtype=False}::test_binary @@ -1190,19 +340,6 @@ tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_para tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_1824_{arg1=-2.0, arg2=array([-3, -2, -1, 1, 2, 3]), dtype=float64, name='fmod', use_dtype=True}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_1825_{arg1=-2.0, arg2=array([-3, -2, -1, 1, 2, 3]), dtype=float64, name='fmod', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_0_{decimals=-2}::test_round_out -tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_1_{decimals=-1}::test_round_out -tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_2_{decimals=0}::test_round -tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_2_{decimals=0}::test_round_out -tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_3_{decimals=1}::test_round_out -tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_4_{decimals=2}::test_round_out -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_negative1 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_positive1 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_negative1 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_positive1 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_negative1 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_positive1 - tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_exp tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_exp2 tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_expm1 @@ -1213,15 +350,6 @@ tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_log2 tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_copysign_combination -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_arccosh -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_arcsinh -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_arctanh -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_cosh -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_sinh -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_tanh - -tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_sign_negative - tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_0_{axis=0, func='nansum', keepdims=True, shape=(2, 3, 4), transpose_axes=True}::test_nansum_all tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_0_{axis=0, func='nansum', keepdims=True, shape=(2, 3, 4), transpose_axes=True}::test_nansum_axis_transposed tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_1_{axis=0, func='nansum', keepdims=True, shape=(2, 3, 4), transpose_axes=False}::test_nansum_all @@ -1280,14 +408,9 @@ tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_3 tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_31_{axis=1, func='nanprod', keepdims=False, shape=(20, 30, 40), transpose_axes=False}::test_nansum_axis_transposed tests/third_party/cupy/math_tests/test_sumprod.py::TestDiff::test_diff_2dim_with_scalar_append -tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_arccos -tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_arcsin -tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_arctan -tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_arctan2 tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_deg2rad tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_hypot tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_rad2deg -tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_tan tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_6_{a_shape=(3, 2), b_shape=(3, 2), shape=(4, 3, 2)}::test_beta tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_7_{a_shape=(3, 2), b_shape=(3, 2), shape=(3, 2)}::test_beta @@ -1460,123 +583,14 @@ tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsWeib tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsWeibull_param_3_{a_shape=(3, 2), shape=(3, 2)}::test_weibull tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsWeibull_param_3_{a_shape=(3, 2), shape=(3, 2)}::test_weibull_for_inf_a -tests/third_party/cupy/sorting_tests/test_search.py::TestSearchSortedWithSorter::test_nonint_sorter - -tests/third_party/cupy/sorting_tests/test_sort.py::TestSort::test_external_sort_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestSort::test_external_sort_contiguous -tests/third_party/cupy/sorting_tests/test_sort.py::TestSort::test_external_sort_invalid_axis1 -tests/third_party/cupy/sorting_tests/test_sort.py::TestSort::test_external_sort_invalid_axis2 -tests/third_party/cupy/sorting_tests/test_sort.py::TestSort::test_external_sort_invalid_negative_axis1 -tests/third_party/cupy/sorting_tests/test_sort.py::TestSort::test_external_sort_invalid_negative_axis2 -tests/third_party/cupy/sorting_tests/test_sort.py::TestSort::test_external_sort_negative_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestSort::test_external_sort_non_contiguous -tests/third_party/cupy/sorting_tests/test_sort.py::TestSort::test_external_sort_none_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestSort::test_external_sort_two_or_more_dim -tests/third_party/cupy/sorting_tests/test_sort.py::TestArgsort_param_1_{external=True}::test_argsort_original_array_not_modified_multi_dim -tests/third_party/cupy/sorting_tests/test_sort.py::TestArgsort_param_1_{external=True}::test_argsort_original_array_not_modified_one_dim -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_axis1 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_axis2 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_negative_axis1 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_negative_axis2 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_negative_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_negative_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_negative_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_non_contiguous -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_sequence_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_zero_dim -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_axis1 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_axis2 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_negative_axis1 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_negative_axis2 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_negative_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_negative_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_negative_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_non_contiguous -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_sequence_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_zero_dim -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_invalid_axis1 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_invalid_axis2 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_invalid_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_invalid_negative_axis1 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_invalid_negative_axis2 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_invalid_negative_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_negative_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_negative_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_non_contiguous -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_none_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_sequence_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_zero_dim -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_invalid_axis1 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_invalid_axis2 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_invalid_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_invalid_negative_axis1 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_invalid_negative_axis2 -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_invalid_negative_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_negative_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_negative_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_non_contiguous -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_none_axis -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_sequence_kth -tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_zero_dim - -tests/third_party/cupy/statistics_tests/test_correlation.py::TestCorrelate_param_0_{mode='valid'}::test_correlate_diff_types -tests/third_party/cupy/statistics_tests/test_correlation.py::TestCorrelate_param_1_{mode='full'}::test_correlate_diff_types -tests/third_party/cupy/statistics_tests/test_correlation.py::TestCorrelate_param_2_{mode='same'}::test_correlate_diff_types -tests/third_party/cupy/statistics_tests/test_correlation.py::TestCorrelate_param_0_{mode='valid'}::test_correlate_diff_types -tests/third_party/cupy/statistics_tests/test_correlation.py::TestCorrelate_param_1_{mode='full'}::test_correlate_diff_types -tests/third_party/cupy/statistics_tests/test_correlation.py::TestCorrelate_param_2_{mode='same'}::test_correlate_diff_types -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedian::test_median_axis1 -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedian::test_median_axis2 -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedian::test_median_invalid_axis -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedian::test_median_keepdims_axis1 -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedian::test_median_keepdims_noaxis -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedian::test_median_noaxis -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedianAxis_param_0_{axis=(0, 1), keepdims=True, shape=(3, 4, 5)}::test_median_axis_sequence -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedianAxis_param_1_{axis=(0, 1), keepdims=False, shape=(3, 4, 5)}::test_median_axis_sequence -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedianAxis_param_2_{axis=(0, -1), keepdims=True, shape=(3, 4, 5)}::test_median_axis_sequence -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedianAxis_param_3_{axis=(0, -1), keepdims=False, shape=(3, 4, 5)}::test_median_axis_sequence -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedianAxis_param_4_{axis=(1, 2), keepdims=True, shape=(3, 4, 5)}::test_median_axis_sequence -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedianAxis_param_5_{axis=(1, 2), keepdims=False, shape=(3, 4, 5)}::test_median_axis_sequence -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedianAxis_param_6_{axis=(1,), keepdims=True, shape=(3, 4, 5)}::test_median_axis_sequence -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMedianAxis_param_7_{axis=(1,), keepdims=False, shape=(3, 4, 5)}::test_median_axis_sequence -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestAverage::test_average_axis -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestAverage::test_average_axis_weights -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_external_mean_all -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_external_mean_axis tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_external_std_all tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_external_std_all_ddof -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_external_std_axis -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_external_std_axis_ddof tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_external_var_all tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_external_var_all_ddof -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_external_var_axis -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_external_var_axis_ddof -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_mean_all -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_mean_axis tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_std_all tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_std_all_ddof -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_std_axis -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_std_axis_ddof tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_var_all tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_var_all_ddof -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_var_axis -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestMeanVar::test_var_axis_ddof -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_0_{func='mean', params=((), None)}::test_external_mean_zero_len -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_1_{func='mean', params=((0,), None)}::test_external_mean_zero_len -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_2_{func='mean', params=((0, 0), None)}::test_external_mean_zero_len -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_3_{func='mean', params=((0, 0), 1)}::test_external_mean_zero_len -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_4_{func='mean', params=((0, 0, 0), None)}::test_external_mean_zero_len -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_5_{func='mean', params=((0, 0, 0), (0, 2))}::test_external_mean_zero_len tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_6_{func='std', params=((), None)}::test_external_mean_zero_len -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_9_{func='std', params=((0, 0), 1)}::test_external_mean_zero_len -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_11_{func='std', params=((0, 0, 0), (0, 2))}::test_external_mean_zero_len tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_12_{func='var', params=((), None)}::test_external_mean_zero_len -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_15_{func='var', params=((0, 0), 1)}::test_external_mean_zero_len -tests/third_party/cupy/statistics_tests/test_meanvar.py::TestProductZeroLength_param_17_{func='var', params=((0, 0, 0), (0, 2))}::test_external_mean_zero_len diff --git a/tests/test_arraycreation.py b/tests/test_arraycreation.py index f8f6007c2f82..c69cdf61dd88 100644 --- a/tests/test_arraycreation.py +++ b/tests/test_arraycreation.py @@ -1,4 +1,3 @@ -import operator import tempfile from math import prod @@ -10,7 +9,6 @@ assert_allclose, assert_almost_equal, assert_array_equal, - assert_raises, ) import dpnp @@ -31,6 +29,7 @@ pytest.param("asarray", {"like": dpnp.array([1, 5])}), pytest.param("ascontiguousarray", {"like": dpnp.zeros(4)}), pytest.param("asfortranarray", {"like": dpnp.empty((2, 4))}), + pytest.param("copy", {"subok": True}), ], ) def test_array_copy_exception(func, kwargs): diff --git a/tests/test_arraymanipulation.py b/tests/test_arraymanipulation.py index 94f54accbc78..1c98c487dfc2 100644 --- a/tests/test_arraymanipulation.py +++ b/tests/test_arraymanipulation.py @@ -13,7 +13,6 @@ from .helper import get_all_dtypes, get_float_complex_dtypes -@pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.parametrize("dtype", get_all_dtypes()) @pytest.mark.parametrize( "data", [[1, 2, 3], [1.0, 2.0, 3.0]], ids=["[1, 2, 3]", "[1., 2., 3.]"] @@ -330,16 +329,31 @@ def test_concatenate_out(self, dtype): assert_array_equal(dp_out.asnumpy(), np_out) assert_array_equal(dp_res.asnumpy(), np_res) + @pytest.mark.usefixtures("allow_fall_back_on_numpy") + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True) + ) + @pytest.mark.parametrize( + "casting", ["no", "equiv", "safe", "same_kind", "unsafe"] + ) + def test_concatenate_casting(self, dtype, casting): + np_a = numpy.arange(2 * 3 * 7, dtype=dtype).reshape((2, 3, 7)) + + dp_a = dpnp.arange(2 * 3 * 7, dtype=dtype).reshape((2, 3, 7)) + + np_res = numpy.concatenate((np_a, np_a), axis=2, casting=casting) + dp_res = dpnp.concatenate((dp_a, dp_a), axis=2, casting=casting) + + assert_array_equal(dp_res.asnumpy(), np_res) + class TestHstack: def test_non_iterable(self): assert_raises(TypeError, dpnp.hstack, 1) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_empty_input(self): - assert_raises(ValueError, dpnp.hstack, ()) + assert_raises(TypeError, dpnp.hstack, ()) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_0D_array(self): b = dpnp.array(2) a = dpnp.array(1) @@ -347,7 +361,6 @@ def test_0D_array(self): desired = dpnp.array([1, 2]) assert_array_equal(res, desired) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_1D_array(self): a = dpnp.array([1]) b = dpnp.array([2]) @@ -355,7 +368,6 @@ def test_1D_array(self): desired = dpnp.array([1, 2]) assert_array_equal(res, desired) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_2D_array(self): a = dpnp.array([[1], [2]]) b = dpnp.array([[1], [2]]) @@ -364,10 +376,15 @@ def test_2D_array(self): assert_array_equal(res, desired) def test_generator(self): - with assert_warns(FutureWarning): - dpnp.hstack((numpy.arange(3) for _ in range(2))) - with assert_warns(FutureWarning): - dpnp.hstack(map(lambda x: x, numpy.ones((3, 2)))) + with pytest.raises(TypeError): + dpnp.hstack((dpnp.arange(3) for _ in range(2))) + with pytest.raises(TypeError): + dpnp.hstack(map(lambda x: x, dpnp.ones((3, 2)))) + + def test_one_element(self): + a = dpnp.array([1]) + res = dpnp.hstack(a) + assert_array_equal(res, a) class TestStack: @@ -605,3 +622,106 @@ def test_2D_array2(self): def test_generator(self): with assert_warns(FutureWarning): dpnp.vstack((numpy.arange(3) for _ in range(2))) + + +class TestAtleast1d: + def test_0D_array(self): + a = dpnp.array(1) + b = dpnp.array(2) + res = [dpnp.atleast_1d(a), dpnp.atleast_1d(b)] + desired = [dpnp.array([1]), dpnp.array([2])] + assert_array_equal(res, desired) + + def test_1D_array(self): + a = dpnp.array([1, 2]) + b = dpnp.array([2, 3]) + res = [dpnp.atleast_1d(a), dpnp.atleast_1d(b)] + desired = [dpnp.array([1, 2]), dpnp.array([2, 3])] + assert_array_equal(res, desired) + + def test_2D_array(self): + a = dpnp.array([[1, 2], [1, 2]]) + b = dpnp.array([[2, 3], [2, 3]]) + res = [dpnp.atleast_1d(a), dpnp.atleast_1d(b)] + desired = [a, b] + assert_array_equal(res, desired) + + def test_3D_array(self): + a = dpnp.array([[1, 2], [1, 2]]) + b = dpnp.array([[2, 3], [2, 3]]) + a = dpnp.array([a, a]) + b = dpnp.array([b, b]) + res = [dpnp.atleast_1d(a), dpnp.atleast_1d(b)] + desired = [a, b] + assert_array_equal(res, desired) + + def test_r1array(self): + assert dpnp.atleast_1d(3).shape == (1,) + assert dpnp.atleast_1d(3j).shape == (1,) + assert dpnp.atleast_1d(3.0).shape == (1,) + assert dpnp.atleast_1d([[2, 3], [4, 5]]).shape == (2, 2) + + +class TestRollaxis: + data = [ + (0, 0), + (0, 1), + (0, 2), + (0, 3), + (0, 4), + (1, 0), + (1, 1), + (1, 2), + (1, 3), + (1, 4), + (2, 0), + (2, 1), + (2, 2), + (2, 3), + (2, 4), + (3, 0), + (3, 1), + (3, 2), + (3, 3), + (3, 4), + ] + + @pytest.mark.parametrize( + ("axis", "start"), + [ + (-5, 0), + (0, -5), + (4, 0), + (0, 5), + ], + ) + def test_exceptions(self, axis, start): + a = dpnp.arange(1 * 2 * 3 * 4).reshape(1, 2, 3, 4) + assert_raises(ValueError, dpnp.rollaxis, a, axis, start) + + def test_results(self): + np_a = numpy.arange(1 * 2 * 3 * 4).reshape(1, 2, 3, 4) + dp_a = dpnp.array(np_a) + for i, j in self.data: + # positive axis, positive start + res = dpnp.rollaxis(dp_a, axis=i, start=j) + exp = numpy.rollaxis(np_a, axis=i, start=j) + assert res.shape == exp.shape + + # negative axis, positive start + ip = i + 1 + res = dpnp.rollaxis(dp_a, axis=-ip, start=j) + exp = numpy.rollaxis(np_a, axis=-ip, start=j) + assert res.shape == exp.shape + + # positive axis, negative start + jp = j + 1 if j < 4 else j + res = dpnp.rollaxis(dp_a, axis=i, start=-jp) + exp = numpy.rollaxis(np_a, axis=i, start=-jp) + assert res.shape == exp.shape + + # negative axis, negative start + ip = i + 1 + jp = j + 1 if j < 4 else j + res = dpnp.rollaxis(dp_a, axis=-ip, start=-jp) + exp = numpy.rollaxis(np_a, axis=-ip, start=-jp) diff --git a/tests/test_bitwise.py b/tests/test_bitwise.py index 9529ac35011c..f8484eaecb5b 100644 --- a/tests/test_bitwise.py +++ b/tests/test_bitwise.py @@ -67,8 +67,6 @@ def test_bitwise_and(self, lhs, rhs, dtype): ) assert_array_equal(dp_a & dp_b, np_a & np_b) - """ - TODO: unmute once dpctl support that if ( not (inp.isscalar(dp_a) or inp.isscalar(dp_b)) and dp_a.shape == dp_b.shape @@ -76,7 +74,6 @@ def test_bitwise_and(self, lhs, rhs, dtype): dp_a &= dp_b np_a &= np_b assert_array_equal(dp_a, np_a) - """ def test_bitwise_or(self, lhs, rhs, dtype): dp_a, dp_b, np_a, np_b = self._test_binary_int( @@ -84,8 +81,6 @@ def test_bitwise_or(self, lhs, rhs, dtype): ) assert_array_equal(dp_a | dp_b, np_a | np_b) - """ - TODO: unmute once dpctl support that if ( not (inp.isscalar(dp_a) or inp.isscalar(dp_b)) and dp_a.shape == dp_b.shape @@ -93,7 +88,6 @@ def test_bitwise_or(self, lhs, rhs, dtype): dp_a |= dp_b np_a |= np_b assert_array_equal(dp_a, np_a) - """ def test_bitwise_xor(self, lhs, rhs, dtype): dp_a, dp_b, np_a, np_b = self._test_binary_int( @@ -101,8 +95,6 @@ def test_bitwise_xor(self, lhs, rhs, dtype): ) assert_array_equal(dp_a ^ dp_b, np_a ^ np_b) - """ - TODO: unmute once dpctl support that if ( not (inp.isscalar(dp_a) or inp.isscalar(dp_b)) and dp_a.shape == dp_b.shape @@ -110,7 +102,6 @@ def test_bitwise_xor(self, lhs, rhs, dtype): dp_a ^= dp_b np_a ^= np_b assert_array_equal(dp_a, np_a) - """ def test_invert(self, lhs, rhs, dtype): dp_a, np_a = self._test_unary_int("invert", lhs, dtype) @@ -122,8 +113,6 @@ def test_left_shift(self, lhs, rhs, dtype): ) assert_array_equal(dp_a << dp_b, np_a << np_b) - """ - TODO: unmute once dpctl support that if ( not (inp.isscalar(dp_a) or inp.isscalar(dp_b)) and dp_a.shape == dp_b.shape @@ -131,7 +120,6 @@ def test_left_shift(self, lhs, rhs, dtype): dp_a <<= dp_b np_a <<= np_b assert_array_equal(dp_a, np_a) - """ def test_right_shift(self, lhs, rhs, dtype): dp_a, dp_b, np_a, np_b = self._test_binary_int( @@ -139,8 +127,6 @@ def test_right_shift(self, lhs, rhs, dtype): ) assert_array_equal(dp_a >> dp_b, np_a >> np_b) - """ - TODO: unmute once dpctl support that if ( not (inp.isscalar(dp_a) or inp.isscalar(dp_b)) and dp_a.shape == dp_b.shape @@ -148,4 +134,3 @@ def test_right_shift(self, lhs, rhs, dtype): dp_a >>= dp_b np_a >>= np_b assert_array_equal(dp_a, np_a) - """ diff --git a/tests/test_copy.py b/tests/test_copy.py new file mode 100644 index 000000000000..e2b161432abd --- /dev/null +++ b/tests/test_copy.py @@ -0,0 +1,74 @@ +import copy + +import numpy +import pytest +from numpy.testing import ( + assert_allclose, + assert_equal, +) + +import dpnp + + +class TestCopyOrder: + a = dpnp.arange(24).reshape(2, 1, 3, 4) + b = a.copy(order="F") + c = dpnp.arange(24).reshape(2, 1, 4, 3).swapaxes(2, 3) + + def check_result(self, x, y, c_contig, f_contig): + assert not (x is y) + assert x.flags.c_contiguous == c_contig + assert x.flags.f_contiguous == f_contig + assert_equal(x, y) + + @pytest.mark.parametrize("arr", [a, b, c]) + def test_order_c(self, arr): + res = arr.copy(order="C") + self.check_result(res, arr, c_contig=True, f_contig=False) + + res = dpnp.copy(arr, order="C") + self.check_result(res, arr, c_contig=True, f_contig=False) + + @pytest.mark.parametrize("arr", [a, b, c]) + def test_order_f(self, arr): + res = arr.copy(order="F") + self.check_result(res, arr, c_contig=False, f_contig=True) + + res = dpnp.copy(arr, order="F") + self.check_result(res, arr, c_contig=False, f_contig=True) + + @pytest.mark.parametrize("arr", [a, b, c]) + def test_order_k(self, arr): + res = arr.copy(order="K") + self.check_result( + res, + arr, + c_contig=arr.flags.c_contiguous, + f_contig=arr.flags.f_contiguous, + ) + + res = dpnp.copy(arr, order="K") + self.check_result( + res, + arr, + c_contig=arr.flags.c_contiguous, + f_contig=arr.flags.f_contiguous, + ) + + res = copy.copy(arr) + self.check_result( + res, + arr, + c_contig=arr.flags.c_contiguous, + f_contig=arr.flags.f_contiguous, + ) + + +@pytest.mark.parametrize( + "val", + [3.7, numpy.arange(7), [2, 7, 3.6], (-3, 4), range(4)], + ids=["scalar", "numpy.array", "list", "tuple", "range"], +) +def test_copy_not_dpnp_array(val): + a = dpnp.copy(val) + assert_allclose(a, val) diff --git a/tests/test_flipping.py b/tests/test_flipping.py new file mode 100644 index 000000000000..36365be1be71 --- /dev/null +++ b/tests/test_flipping.py @@ -0,0 +1,169 @@ +from math import prod + +import numpy +import pytest +from numpy.testing import ( + assert_equal, +) + +import dpnp + +from .helper import ( + get_all_dtypes, +) + + +class TestFlip: + @pytest.mark.parametrize("dtype", get_all_dtypes()) + def test_arange_2d_default_axis(self, dtype): + sh = (2, 3) if dtype != dpnp.bool else (1, 1) + dp_a = dpnp.arange(prod(sh), dtype=dtype).reshape(sh) + np_a = numpy.arange(prod(sh), dtype=dtype).reshape(sh) + + assert_equal(dpnp.flip(dp_a), numpy.flip(np_a)) + + @pytest.mark.parametrize("axis", list(range(3))) + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True) + ) + def test_arange_3d(self, axis, dtype): + sh = (2, 2, 2) + dp_a = dpnp.arange(prod(sh), dtype=dtype).reshape(sh) + np_a = numpy.arange(prod(sh), dtype=dtype).reshape(sh) + + assert_equal(dpnp.flip(dp_a, axis=axis), numpy.flip(np_a, axis=axis)) + + @pytest.mark.parametrize("axis", [(), (0, 2), (1, 2)]) + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True) + ) + def test_arange_3d_multiple_axes(self, axis, dtype): + sh = (2, 2, 2) + dp_a = dpnp.arange(prod(sh), dtype=dtype).reshape(sh) + np_a = numpy.arange(prod(sh), dtype=dtype).reshape(sh) + + assert_equal(dpnp.flip(dp_a, axis=axis), numpy.flip(np_a, axis=axis)) + + @pytest.mark.parametrize("axis", list(range(4))) + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True) + ) + def test_arange_4d(self, axis, dtype): + sh = (2, 3, 4, 5) + dp_a = dpnp.arange(prod(sh), dtype=dtype).reshape(sh) + np_a = numpy.arange(prod(sh), dtype=dtype).reshape(sh) + + assert_equal(dpnp.flip(dp_a, axis=axis), numpy.flip(np_a, axis=axis)) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True) + ) + def test_lr_equivalent(self, dtype): + dp_a = dpnp.arange(4, dtype=dtype) + dp_a = dp_a[:, dpnp.newaxis] + dp_a[dpnp.newaxis, :] + assert_equal(dpnp.flip(dp_a, 1), dpnp.fliplr(dp_a)) + + np_a = numpy.arange(4, dtype=dtype) + np_a = numpy.add.outer(np_a, np_a) + assert_equal(dpnp.flip(dp_a, 1), numpy.flip(np_a, 1)) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True) + ) + def test_ud_equivalent(self, dtype): + dp_a = dpnp.arange(4, dtype=dtype) + dp_a = dp_a[:, dpnp.newaxis] + dp_a[dpnp.newaxis, :] + assert_equal(dpnp.flip(dp_a, 0), dpnp.flipud(dp_a)) + + np_a = numpy.arange(4, dtype=dtype) + np_a = numpy.add.outer(np_a, np_a) + assert_equal(dpnp.flip(dp_a, 0), numpy.flip(np_a, 0)) + + @pytest.mark.parametrize( + "x, axis", + [ + pytest.param(dpnp.ones(4), 1, id="1-d, axis=1"), + pytest.param(dpnp.ones((4, 4)), 2, id="2-d, axis=2"), + pytest.param(dpnp.ones((4, 4)), -3, id="2-d, axis=-3"), + pytest.param(dpnp.ones((4, 4)), (0, 3), id="2-d, axis=(0, 3)"), + ], + ) + def test_axes(self, x, axis): + with pytest.raises(numpy.AxisError): + dpnp.flip(x, axis=axis) + + +class TestFliplr: + @pytest.mark.parametrize("dtype", get_all_dtypes()) + def test_arange(self, dtype): + sh = (2, 3) if dtype != dpnp.bool else (1, 1) + dp_a = dpnp.arange(prod(sh), dtype=dtype).reshape(sh) + np_a = numpy.arange(prod(sh), dtype=dtype).reshape(sh) + + assert_equal(dpnp.fliplr(dp_a), numpy.fliplr(np_a)) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True) + ) + def test_equivalent(self, dtype): + dp_a = dpnp.arange(4, dtype=dtype) + dp_a = dp_a[:, dpnp.newaxis] + dp_a[dpnp.newaxis, :] + assert_equal(dpnp.fliplr(dp_a), dp_a[:, ::-1]) + + np_a = numpy.arange(4, dtype=dtype) + np_a = numpy.add.outer(np_a, np_a) + assert_equal(dpnp.fliplr(dp_a), numpy.fliplr(np_a)) + + @pytest.mark.parametrize( + "val", + [-1.2, numpy.arange(7), [2, 7, 3.6], (-3, 4), range(4)], + ids=["scalar", "numpy.array", "list", "tuple", "range"], + ) + def test_raises_array_type(self, val): + with pytest.raises( + TypeError, match="An array must be any of supported type, but got" + ): + dpnp.fliplr(val) + + def test_raises_1d(self): + a = dpnp.ones(4) + with pytest.raises(ValueError, match="Input must be >= 2-d, but got"): + dpnp.fliplr(a) + + +class TestFlipud: + @pytest.mark.parametrize("dtype", get_all_dtypes()) + def test_arange(self, dtype): + sh = (2, 3) if dtype != dpnp.bool else (1, 1) + dp_a = dpnp.arange(prod(sh), dtype=dtype).reshape(sh) + np_a = numpy.arange(prod(sh), dtype=dtype).reshape(sh) + + assert_equal(dpnp.flipud(dp_a), numpy.flipud(np_a)) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True) + ) + def test_equivalent(self, dtype): + dp_a = dpnp.arange(4, dtype=dtype) + dp_a = dp_a[:, dpnp.newaxis] + dp_a[dpnp.newaxis, :] + assert_equal(dpnp.flipud(dp_a), dp_a[::-1, :]) + + np_a = numpy.arange(4, dtype=dtype) + np_a = numpy.add.outer(np_a, np_a) + assert_equal(dpnp.flipud(dp_a), numpy.flipud(np_a)) + + @pytest.mark.parametrize( + "val", + [3.4, numpy.arange(6), [2, -1.7, 6], (-2, 4), range(5)], + ids=["scalar", "numpy.array", "list", "tuple", "range"], + ) + def test_raises_array_type(self, val): + with pytest.raises( + TypeError, match="An array must be any of supported type, but got" + ): + dpnp.flipud(val) + + def test_raises_0d(self): + a = dpnp.array(3) + with pytest.raises(ValueError, match="Input must be >= 1-d, but got"): + dpnp.flipud(a) diff --git a/tests/test_indexing.py b/tests/test_indexing.py index 67600264356d..03541dc2d55d 100644 --- a/tests/test_indexing.py +++ b/tests/test_indexing.py @@ -341,76 +341,124 @@ def test_place3(arr, mask, vals): assert_array_equal(a, ia) -@pytest.mark.parametrize("v", [0, 1, 2, 3, 4], ids=["0", "1", "2", "3", "4"]) -@pytest.mark.parametrize("ind", [0, 1, 2, 3], ids=["0", "1", "2", "3"]) +@pytest.mark.parametrize("array_dtype", get_all_dtypes()) @pytest.mark.parametrize( - "array", - [ - [[0, 0], [0, 0]], - [[1, 2], [1, 2]], - [[1, 2], [3, 4]], - [[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]], - [ - [[[1, 2], [3, 4]], [[1, 2], [2, 1]]], - [[[1, 3], [3, 1]], [[0, 1], [1, 3]]], - ], - ], - ids=[ - "[[0, 0], [0, 0]]", - "[[1, 2], [1, 2]]", - "[[1, 2], [3, 4]]", - "[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]", - "[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]", - ], + "indices_dtype", [dpnp.int32, dpnp.int64], ids=["int32", "int64"] ) -def test_put(array, ind, v): - a = numpy.array(array) +@pytest.mark.parametrize( + "indices", [[-2, 2], [-5, 4]], ids=["[-2, 2]", "[-5, 4]"] +) +@pytest.mark.parametrize( + "vals", + [0, [1, 2], (2, 2), dpnp.array([1, 2])], + ids=["0", "[1, 2]", "(2, 2)", "dpnp.array([1,2])"], +) +@pytest.mark.parametrize("mode", ["clip", "wrap"], ids=["clip", "wrap"]) +def test_put_1d(indices, vals, array_dtype, indices_dtype, mode): + a = numpy.array([-2, -1, 0, 1, 2], dtype=array_dtype) + b = numpy.copy(a) ia = dpnp.array(a) - numpy.put(a, ind, v) - dpnp.put(ia, ind, v) + ib = dpnp.array(b) + ind = numpy.array(indices, dtype=indices_dtype) + iind = dpnp.array(ind) + + # TODO: remove when #1382(dpctl) is solved + if dpnp.is_supported_array_type(vals): + vals = dpnp.astype(vals, ia.dtype) + + numpy.put(a, ind, vals, mode=mode) + dpnp.put(ia, iind, vals, mode=mode) assert_array_equal(a, ia) + b.put(ind, vals, mode=mode) + ib.put(iind, vals, mode=mode) + assert_array_equal(b, ib) + +@pytest.mark.parametrize("array_dtype", get_all_dtypes()) @pytest.mark.parametrize( - "v", [[10, 20], [30, 40]], ids=["[10, 20]", "[30, 40]"] + "indices_dtype", [dpnp.int32, dpnp.int64], ids=["int32", "int64"] ) -@pytest.mark.parametrize("ind", [[0, 1], [2, 3]], ids=["[0, 1]", "[2, 3]"]) +@pytest.mark.parametrize("vals", [[10, 20]], ids=["[10, 20]"]) @pytest.mark.parametrize( - "array", + "indices", [ - [[0, 0], [0, 0]], - [[1, 2], [1, 2]], - [[1, 2], [3, 4]], - [[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]], - [ - [[[1, 2], [3, 4]], [[1, 2], [2, 1]]], - [[[1, 3], [3, 1]], [[0, 1], [1, 3]]], - ], + [0, 7], + [3, 4], + [-9, 8], ], ids=[ - "[[0, 0], [0, 0]]", - "[[1, 2], [1, 2]]", - "[[1, 2], [3, 4]]", - "[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]", - "[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]", + "[0, 7]", + "[3, 4]", + "[-9, 8]", ], ) -def test_put2(array, ind, v): - a = numpy.array(array) +@pytest.mark.parametrize("mode", ["clip", "wrap"], ids=["clip", "wrap"]) +def test_put_2d(array_dtype, indices_dtype, indices, vals, mode): + a = numpy.array([[-1, 0, 1], [-2, -3, -4], [2, 3, 4]], dtype=array_dtype) ia = dpnp.array(a) - numpy.put(a, ind, v) - dpnp.put(ia, ind, v) + ind = numpy.array(indices, dtype=indices_dtype) + iind = dpnp.array(ind) + numpy.put(a, ind, vals, mode=mode) + dpnp.put(ia, iind, vals, mode=mode) assert_array_equal(a, ia) -def test_put3(): +@pytest.mark.usefixtures("allow_fall_back_on_numpy") +def test_put_2d_ind(): a = numpy.arange(5) ia = dpnp.array(a) - dpnp.put(ia, [0, 2], [-44, -55]) - numpy.put(a, [0, 2], [-44, -55]) + ind = numpy.array([[3, 0, 2, 1]]) + iind = dpnp.array(ind) + numpy.put(a, ind, 10) + dpnp.put(ia, iind, 10) assert_array_equal(a, ia) +@pytest.mark.parametrize( + "shape", + [ + (0,), + (3,), + (4,), + ], + ids=[ + "(0,)", + "(3,)", + "(4,)", + ], +) +@pytest.mark.parametrize("mode", ["clip", "wrap"], ids=["clip", "wrap"]) +def test_put_invalid_shape(shape, mode): + a = dpnp.arange(7) + ind = dpnp.array([2]) + vals = dpnp.ones(shape, dtype=a.dtype) + # vals must be broadcastable to the shape of ind` + with pytest.raises(ValueError): + dpnp.put(a, ind, vals, mode=mode) + + +@pytest.mark.parametrize( + "axis", + [ + 1.0, + (0,), + [0, 1], + ], + ids=[ + "1.0", + "(0,)", + "[0, 1]", + ], +) +def test_put_invalid_axis(axis): + a = dpnp.arange(6).reshape(2, 3) + ind = dpnp.array([1]) + vals = [1] + with pytest.raises(TypeError): + dpnp.put(a, ind, vals, axis=axis) + + @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_put_along_axis_val_int(): a = numpy.arange(16).reshape(4, 4) diff --git a/tests/test_logic.py b/tests/test_logic.py index 7be9e6e1ac8c..b9d2a9b4303a 100644 --- a/tests/test_logic.py +++ b/tests/test_logic.py @@ -44,16 +44,13 @@ def test_all(type, shape): assert_allclose(dpnp_res, np_res) -@pytest.mark.skipif( - not has_support_aspect64(), reason="Aborted on Iris Xe: SAT-5988" -) -@pytest.mark.parametrize("type", get_all_dtypes(no_bool=True, no_complex=True)) -def test_allclose(type): +@pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) +def test_allclose(dtype): a = numpy.random.rand(10) b = a + numpy.random.rand(10) * 1e-8 - dpnp_a = dpnp.array(a, dtype=type) - dpnp_b = dpnp.array(b, dtype=type) + dpnp_a = dpnp.array(a, dtype=dtype) + dpnp_b = dpnp.array(b, dtype=dtype) np_res = numpy.allclose(a, b) dpnp_res = dpnp.allclose(dpnp_a, dpnp_b) @@ -68,6 +65,63 @@ def test_allclose(type): assert_allclose(dpnp_res, np_res) +class TestAllClose: + @pytest.mark.parametrize("val", [1.0, 3, numpy.inf, -numpy.inf, numpy.nan]) + def test_input_0d(self, val): + dp_arr = dpnp.array(val) + np_arr = numpy.array(val) + + # array & scalar + dp_res = dpnp.allclose(dp_arr, val) + np_res = numpy.allclose(np_arr, val) + assert_allclose(dp_res, np_res) + + # scalar & array + dp_res = dpnp.allclose(val, dp_arr) + np_res = numpy.allclose(val, np_arr) + assert_allclose(dp_res, np_res) + + # two arrays + dp_res = dpnp.allclose(dp_arr, dp_arr) + np_res = numpy.allclose(np_arr, np_arr) + assert_allclose(dp_res, np_res) + + @pytest.mark.parametrize("sh_a", [(10,), (10, 10)]) + @pytest.mark.parametrize("sh_b", [(1, 10), (1, 10, 1)]) + def test_broadcast(self, sh_a, sh_b): + dp_a = dpnp.ones(sh_a) + dp_b = dpnp.ones(sh_b) + + np_a = numpy.ones(sh_a) + np_b = numpy.ones(sh_b) + + dp_res = dpnp.allclose(dp_a, dp_b) + np_res = numpy.allclose(np_a, np_b) + assert_allclose(dp_res, np_res) + + def test_input_as_scalars(self): + with pytest.raises(NotImplementedError): + dpnp.allclose(1.0, 1.0) + + @pytest.mark.parametrize("val", [[1.0], (-3, 7), numpy.arange(5)]) + def test_wrong_input_arrays(self, val): + with pytest.raises(NotImplementedError): + dpnp.allclose(val, val) + + @pytest.mark.parametrize( + "tol", [[0.001], (1.0e-6,), dpnp.array(1.0e-3), numpy.array([1.0e-5])] + ) + def test_wrong_tols(self, tol): + a = dpnp.ones(10) + b = dpnp.ones(10) + + for kw in [{"rtol": tol}, {"atol": tol}, {"rtol": tol, "atol": tol}]: + with pytest.raises( + TypeError, match=r"An argument .* must be a scalar, but got" + ): + dpnp.allclose(a, b, **kw) + + @pytest.mark.parametrize("type", get_all_dtypes()) @pytest.mark.parametrize( "shape", diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 2b5c98083e29..adb556fb79c7 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -13,8 +13,9 @@ import dpnp from .helper import ( + assert_dtype_allclose, get_all_dtypes, - get_float_complex_dtypes, + get_complex_dtypes, get_float_dtypes, has_support_aspect64, is_cpu_device, @@ -119,7 +120,7 @@ def array_or_scalar(xp, data, dtype=None): return xp.array(data, dtype=dtype) - def _test_mathematical(self, name, dtype, lhs, rhs): + def _test_mathematical(self, name, dtype, lhs, rhs, check_type=True): a_dpnp = self.array_or_scalar(dpnp, lhs, dtype=dtype) b_dpnp = self.array_or_scalar(dpnp, rhs, dtype=dtype) @@ -137,22 +138,13 @@ def _test_mathematical(self, name, dtype, lhs, rhs): else: result = getattr(dpnp, name)(a_dpnp, b_dpnp) expected = getattr(numpy, name)(a_np, b_np) - if ( - name == "remainder" - and result.dtype != expected.dtype - and not has_support_aspect64() - ): - pytest.skip("skipping since output is promoted differently") - assert_allclose(result, expected, rtol=1e-6) + assert_dtype_allclose(result, expected, check_type) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_add(self, dtype, lhs, rhs): - self._test_mathematical("add", dtype, lhs, rhs) + self._test_mathematical("add", dtype, lhs, rhs, check_type=False) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") - @pytest.mark.parametrize( - "dtype", get_all_dtypes(no_bool=True, no_complex=True) - ) + @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) def test_arctan2(self, dtype, lhs, rhs): self._test_mathematical("arctan2", dtype, lhs, rhs) @@ -172,17 +164,18 @@ def test_divide(self, dtype, lhs, rhs): "dtype", get_all_dtypes(no_bool=True, no_complex=True) ) def test_fmod(self, dtype, lhs, rhs): - if dtype == None and rhs == 0.3 and not has_support_aspect64(): + if rhs == 0.3: """ - Due to accuracy reason NumPy behaves differently, when: + Due to accuracy reason, the results are different for `float32` and `float64` >>> numpy.fmod(numpy.array([3.9], dtype=numpy.float32), 0.3) array([0.29999995], dtype=float32) - while numpy with float64 returns something around zero which is aligned with dpnp: + >>> numpy.fmod(numpy.array([3.9], dtype=numpy.float64), 0.3) array([9.53674318e-08]) + On a gpu without support for `float64`, dpnp produces results similar to the second one. """ - pytest.skip("missaligned between numpy results") - self._test_mathematical("fmod", dtype, lhs, rhs) + pytest.skip("Due to accuracy reason, the results are different.") + self._test_mathematical("fmod", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) def test_floor_divide(self, dtype, lhs, rhs): @@ -190,7 +183,9 @@ def test_floor_divide(self, dtype, lhs, rhs): pytest.skip( "In this case, a different result, but similar to xp.floor(xp.divide(lhs, rhs)." ) - self._test_mathematical("floor_divide", dtype, lhs, rhs) + self._test_mathematical( + "floor_divide", dtype, lhs, rhs, check_type=False + ) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.parametrize( @@ -204,30 +199,45 @@ def test_hypot(self, dtype, lhs, rhs): "dtype", get_all_dtypes(no_bool=True, no_complex=True) ) def test_maximum(self, dtype, lhs, rhs): - self._test_mathematical("maximum", dtype, lhs, rhs) + self._test_mathematical("maximum", dtype, lhs, rhs, check_type=False) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True) ) def test_minimum(self, dtype, lhs, rhs): - self._test_mathematical("minimum", dtype, lhs, rhs) + self._test_mathematical("minimum", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_multiply(self, dtype, lhs, rhs): - self._test_mathematical("multiply", dtype, lhs, rhs) + self._test_mathematical("multiply", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) def test_remainder(self, dtype, lhs, rhs): - self._test_mathematical("remainder", dtype, lhs, rhs) + if ( + dtype in [dpnp.int32, dpnp.int64, None] + and rhs == 0.3 + and not has_support_aspect64() + ): + """ + Due to accuracy reason, the results are different for `float32` and `float64` + >>> numpy.remainder(numpy.array([6, 3], dtype='i4'), 0.3, dtype='f8') + array([2.22044605e-16, 1.11022302e-16]) + + >>> numpy.remainder(numpy.array([6, 3], dtype='i4'), 0.3, dtype='f4') + usm_ndarray([0.29999977, 0.2999999 ], dtype=float32) + On a gpu without support for `float64`, dpnp produces results similar to the second one. + """ + pytest.skip("Due to accuracy reason, the results are different.") + self._test_mathematical("remainder", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_power(self, dtype, lhs, rhs): - self._test_mathematical("power", dtype, lhs, rhs) + self._test_mathematical("power", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_subtract(self, dtype, lhs, rhs): - self._test_mathematical("subtract", dtype, lhs, rhs) + self._test_mathematical("subtract", dtype, lhs, rhs, check_type=False) @pytest.mark.usefixtures("suppress_divide_invalid_numpy_warnings") @@ -265,14 +275,23 @@ def test_op_with_scalar(array, val, func, data_type, val_type): val_ = val_type(val) if func == "power": - if val_ == 0 and numpy.issubdtype(data_type, numpy.complexfloating): + if ( + val_ == 0 + and numpy.issubdtype(data_type, numpy.complexfloating) + and not dpnp.all(dpnp_a) + ): pytest.skip( "(0j ** 0) is different: (NaN + NaNj) in dpnp and (1 + 0j) in numpy" ) - elif is_cpu_device() and data_type == dpnp.complex128: - # TODO: discuss the bahavior with OneMKL team + # TODO: Remove when #1378 (dpctl) is solved + elif ( + is_cpu_device() + and dpnp_a.dtype == dpnp.complex128 + and dpnp_a.size >= 8 + and not dpnp.all(dpnp_a) + ): pytest.skip( - "(0j ** 5) is different: (NaN + NaNj) in dpnp and (0j) in numpy" + "[..., 0j ** val] is different for x.size >= 8: [..., NaN + NaNj] in dpnp and [..., 0 + 0j] in numpy" ) if func == "subtract" and val_type == bool and data_type == dpnp.bool: @@ -432,6 +451,48 @@ def test_sign_boolean(): dpnp.sign(dpnp_a) +@pytest.mark.parametrize( + "data", + [[2, 0, -2], [1.1, -1.1]], + ids=["[2, 0, -2]", "[1.1, -1.1]"], +) +@pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) +def test_signbit(data, dtype): + np_a = numpy.array(data, dtype=dtype) + dpnp_a = dpnp.array(data, dtype=dtype) + + result = dpnp.signbit(dpnp_a) + expected = numpy.signbit(np_a) + assert_allclose(result, expected) + + +@pytest.mark.parametrize("dtype", get_complex_dtypes()) +def test_projection_infinity(dtype): + X = [ + complex(1, 2), + complex(dpnp.inf, -1), + complex(0, -dpnp.inf), + complex(-dpnp.inf, dpnp.nan), + ] + Y = [ + complex(1, 2), + complex(dpnp.inf, -0.0), + complex(dpnp.inf, -0.0), + complex(dpnp.inf, 0.0), + ] + + result = dpnp.proj(dpnp.array(X, dtype=dtype)) + expected = dpnp.array(Y, dtype=dtype) + assert_allclose(result, expected) + + +@pytest.mark.parametrize("dtype", get_all_dtypes()) +def test_projection(dtype): + result = dpnp.proj(dpnp.array(1, dtype=dtype)) + expected = dpnp.array(complex(1, 0)) + assert_allclose(result, expected) + + @pytest.mark.parametrize("val_type", get_all_dtypes(no_none=True)) @pytest.mark.parametrize("data_type", get_all_dtypes()) @pytest.mark.parametrize("val", [1.5, 1, 5], ids=["1.5", "1", "5"]) @@ -460,10 +521,17 @@ def test_power(array, val, data_type, val_type): dpnp_a = dpnp.array(array, dtype=data_type) val_ = val_type(val) - if is_cpu_device() and dpnp.complex128 in (data_type, val_type): - # TODO: discuss the behavior with OneMKL team + # TODO: Remove when #1378 (dpctl) is solved + if ( + is_cpu_device() + and ( + dpnp.complex128 in (data_type, val_type) + or dpnp.complex64 in (data_type, val_type) + ) + and dpnp_a.size >= 8 + ): pytest.skip( - "(0j ** 5) is different: (NaN + NaNj) in dpnp and (0j) in numpy" + "[..., 0j ** val] is different for x.size >= 8: [..., NaN + NaNj] in dpnp and [..., 0 + 0j] in numpy" ) result = dpnp.power(dpnp_a, val_) @@ -938,7 +1006,9 @@ def test_invalid_out(self, out): class TestPower: - @pytest.mark.parametrize("dtype", get_float_complex_dtypes()) + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True) + ) def test_power(self, dtype): array1_data = numpy.arange(10) array2_data = numpy.arange(5, 15) @@ -955,11 +1025,9 @@ def test_power(self, dtype): np_array2 = numpy.array(array2_data, dtype=dtype) expected = numpy.power(np_array1, np_array2, out=out) - assert_allclose(expected, result) + assert_allclose(expected, result, rtol=1e-06) - @pytest.mark.parametrize( - "dtype", get_all_dtypes(no_complex=True, no_none=True) - ) + @pytest.mark.parametrize("dtype", get_all_dtypes(no_none=True)) def test_out_dtypes(self, dtype): size = 2 if dtype == dpnp.bool else 5 @@ -971,27 +1039,37 @@ def test_out_dtypes(self, dtype): dp_array1 = dpnp.arange(size, 2 * size, dtype=dtype) dp_array2 = dpnp.arange(size, dtype=dtype) dp_out = dpnp.empty(size, dtype=dpnp.complex64) - result = dpnp.power(dp_array1, dp_array2, out=dp_out) + if dtype != dpnp.complex64: + # dtype of out mismatches types of input arrays + with pytest.raises(TypeError): + dpnp.power(dp_array1, dp_array2, out=dp_out) - assert_array_equal(expected, result) + # allocate new out with expected type + if dtype == dpnp.bool: + out_dtype = numpy.int8 + else: + out_dtype = dtype + dp_out = dpnp.empty(size, dtype=out_dtype) - @pytest.mark.parametrize( - "dtype", get_all_dtypes(no_bool=True, no_complex=True, no_none=True) - ) + result = dpnp.power(dp_array1, dp_array2, out=dp_out) + assert_allclose(expected, result, rtol=1e-06) + + @pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True)) def test_out_overlap(self, dtype): - size = 5 + size = 10 + # DPNP + dp_a = dpnp.arange(2 * size, dtype=dtype) + dpnp.power(dp_a[size::], dp_a[::2], out=dp_a[:size:]), + # original np_a = numpy.arange(2 * size, dtype=dtype) - expected = numpy.power(np_a[size::], np_a[::2], out=np_a[:size:]) + numpy.power(np_a[size::], np_a[::2], out=np_a[:size:]) - dp_a = dpnp.arange(2 * size, dtype=dtype) - result = dpnp.power(dp_a[size::], dp_a[::2], out=dp_a[:size:]) - - assert_allclose(expected, result) - assert_allclose(dp_a, np_a) + rtol = 1e-05 if dtype is dpnp.complex64 else 1e-07 + assert_allclose(np_a, dp_a, rtol=rtol) @pytest.mark.parametrize( - "dtype", get_all_dtypes(no_bool=True, no_complex=True, no_none=True) + "dtype", get_all_dtypes(no_bool=True, no_none=True) ) def test_inplace_strided_out(self, dtype): size = 5 @@ -1008,9 +1086,9 @@ def test_inplace_strided_out(self, dtype): "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) def test_invalid_shape(self, shape): - dp_array1 = dpnp.arange(10, dtype=dpnp.float64) - dp_array2 = dpnp.arange(5, 15, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + dp_array1 = dpnp.arange(10, dtype=dpnp.float32) + dp_array2 = dpnp.arange(5, 15, dtype=dpnp.float32) + dp_out = dpnp.empty(shape, dtype=dpnp.float32) with pytest.raises(ValueError): dpnp.power(dp_array1, dp_array2, out=dp_out) @@ -1045,19 +1123,24 @@ def test_complex_values(self): def test_integer_power_of_0_or_1(self, val, dtype): np_arr = numpy.arange(10, dtype=dtype) dp_arr = dpnp.array(np_arr) - func = lambda x: 1**x + func = lambda x: val**x assert_equal(func(np_arr), func(dp_arr)) @pytest.mark.parametrize("dtype", [dpnp.int32, dpnp.int64]) def test_integer_to_negative_power(self, dtype): - ones = dpnp.ones(10, dtype=dtype) a = dpnp.arange(2, 10, dtype=dtype) - b = dpnp.full(10, -2, dtype=dtype) - - assert_array_equal(ones ** (-2), ones) - assert_equal(a ** (-3), 0) # positive integer to negative integer power - assert_equal(b ** (-4), 0) # negative integer to negative integer power + b = dpnp.full(8, -2, dtype=dtype) + zeros = dpnp.zeros(8, dtype=dtype) + ones = dpnp.ones(8, dtype=dtype) + + assert_array_equal(ones ** (-2), zeros) + assert_equal( + a ** (-3), zeros + ) # positive integer to negative integer power + assert_equal( + b ** (-4), zeros + ) # negative integer to negative integer power def test_float_to_inf(self): a = numpy.array( @@ -1177,3 +1260,31 @@ def test_mean_scalar(self): result = dp_array.mean() expected = np_array.mean() assert_allclose(expected, result) + + +@pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True, no_complex=True) +) +def test_inplace_remainder(dtype): + size = 21 + np_a = numpy.arange(size, dtype=dtype) + dp_a = dpnp.arange(size, dtype=dtype) + + np_a %= 4 + dp_a %= 4 + + assert_allclose(dp_a, np_a) + + +@pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_none=True, no_complex=True) +) +def test_inplace_floor_divide(dtype): + size = 21 + np_a = numpy.arange(size, dtype=dtype) + dp_a = dpnp.arange(size, dtype=dtype) + + np_a //= 4 + dp_a //= 4 + + assert_allclose(dp_a, np_a) diff --git a/tests/test_strides.py b/tests/test_strides.py index e39414a7bd7f..1ddfcd14eb5a 100644 --- a/tests/test_strides.py +++ b/tests/test_strides.py @@ -74,6 +74,7 @@ def test_strides(func_name, dtype): "sinh", "sqrt", "square", + "tan", "tanh", "trunc", ], @@ -93,7 +94,7 @@ def test_strides_1arg(func_name, dtype, shape): numpy_func = _getattr(numpy, func_name) expected = numpy_func(b) - assert_allclose(result, expected) + assert_allclose(result, expected, rtol=1e-06) @pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) @@ -130,21 +131,6 @@ def test_strides_reciprocal(dtype, shape): assert_allclose(result, expected, rtol=1e-06) -@pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) -@pytest.mark.parametrize("shape", [(10,)], ids=["(10,)"]) -def test_strides_tan(dtype, shape): - a = numpy.arange(numpy.prod(shape), dtype=dtype).reshape(shape) - b = a[::2] - - dpa = dpnp.reshape(dpnp.arange(numpy.prod(shape), dtype=dtype), shape) - dpb = dpa[::2] - - result = dpnp.tan(dpb) - expected = numpy.tan(b) - - assert_allclose(result, expected, rtol=1e-06) - - @pytest.mark.parametrize( "func_name", [ @@ -161,6 +147,7 @@ def test_strides_tan(dtype, shape): ) @pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) @pytest.mark.parametrize("shape", [(3, 3)], ids=["(3, 3)"]) +@pytest.mark.usefixtures("suppress_invalid_numpy_warnings") def test_strides_2args(func_name, dtype, shape): a = numpy.arange(numpy.prod(shape), dtype=dtype).reshape(shape) b = a.T @@ -277,8 +264,8 @@ def test_strided_out_2args(func_name, dtype): np_res = _getattr(numpy, func_name)(np_a, np_b, out=np_out) dp_res = _getattr(dpnp, func_name)(dp_a, dp_b, out=dp_out) - assert_allclose(dp_res.asnumpy(), np_res) - assert_allclose(dp_out.asnumpy(), np_out) + assert_allclose(dp_res.asnumpy(), np_res, rtol=1e-06) + assert_allclose(dp_out.asnumpy(), np_out, rtol=1e-06) @pytest.mark.parametrize("func_name", ["add", "multiply", "power", "subtract"]) @@ -328,7 +315,6 @@ def test_strided_in_out_2args_diff_out_dtype(func_name, dtype): @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True, no_none=True) ) -@pytest.mark.skip("dpctl doesn't support overlap of arrays") def test_strided_in_2args_overlap(func_name, dtype): size = 5 @@ -350,7 +336,6 @@ def test_strided_in_2args_overlap(func_name, dtype): @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True, no_none=True) ) -@pytest.mark.skip("dpctl doesn't support overlap of arrays") def test_strided_in_out_2args_overlap(func_name, dtype): sh = (4, 3, 2) prod = numpy.prod(sh) diff --git a/tests/test_sum.py b/tests/test_sum.py index 8701beafc8a3..16b17847f270 100644 --- a/tests/test_sum.py +++ b/tests/test_sum.py @@ -1,12 +1,19 @@ import numpy import pytest +from numpy.testing import ( + assert_array_equal, +) import dpnp -from tests.helper import get_float_dtypes, has_support_aspect64 +from tests.helper import ( + assert_dtype_allclose, + get_float_dtypes, + has_support_aspect64, +) -# Note: numpy.sum() always upcast integers to (u)int64 and float32 to -# float64 for dtype=None. `np.sum` does that too for integers, but not for +# Note: dpnp.sum() always upcast integers to (u)int64 and float32 to +# float64 for dtype=None. `numpy.sum()` does that too for integers, but not for # float32, so we need to special-case it for these tests @pytest.mark.parametrize("dtype", get_float_dtypes()) def test_sum_float(dtype): @@ -22,11 +29,8 @@ def test_sum_float(dtype): for axis in range(len(a)): result = dpnp.sum(ia, axis=axis) - if dtype == dpnp.float32 and has_support_aspect64(): - expected = numpy.sum(a, axis=axis, dtype=numpy.float64) - else: - expected = numpy.sum(a, axis=axis) - numpy.testing.assert_array_equal(expected, result) + expected = numpy.sum(a, axis=axis) + assert_dtype_allclose(result, expected) def test_sum_int(): @@ -35,7 +39,7 @@ def test_sum_int(): result = dpnp.sum(ia) expected = numpy.sum(a) - numpy.testing.assert_array_equal(expected, result) + assert_array_equal(expected, result) def test_sum_axis(): @@ -54,4 +58,4 @@ def test_sum_axis(): expected = numpy.sum(a, axis=1, dtype=numpy.float64) else: expected = numpy.sum(a, axis=1) - numpy.testing.assert_array_equal(expected, result) + assert_array_equal(expected, result) diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py index 78f98c61581c..f5b0248f33c7 100644 --- a/tests/test_sycl_queue.py +++ b/tests/test_sycl_queue.py @@ -94,13 +94,13 @@ def vvsort(val, vec, size, xp): ids=[device.filter_string for device in valid_devices], ) def test_array_creation(func, arg, kwargs, device): - numpy_array = getattr(numpy, func)(*arg, **kwargs) - dpnp_kwargs = dict(kwargs) dpnp_kwargs["device"] = device dpnp_array = getattr(dpnp, func)(*arg, **dpnp_kwargs) - assert_allclose(numpy_array, dpnp_array) + numpy_array = getattr(numpy, func)(*arg, dtype=dpnp_array.dtype, **kwargs) + + assert_dtype_allclose(dpnp_array, numpy_array) assert dpnp_array.sycl_device == device @@ -225,8 +225,18 @@ def test_meshgrid(device_x, device_y): "func,data", [ pytest.param("abs", [-1.2, 1.2]), + pytest.param("arccos", [-0.5, 0.0, 0.5]), + pytest.param("arccosh", [1.5, 3.5, 5.0]), + pytest.param("arcsin", [-0.5, 0.0, 0.5]), + pytest.param("arcsinh", [-5.0, -3.5, 0.0, 3.5, 5.0]), + pytest.param("arctan", [-1.0, 0.0, 1.0]), + pytest.param("arctanh", [-0.5, 0.0, 0.5]), pytest.param("ceil", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), pytest.param("conjugate", [[1.0 + 1.0j, 0.0], [0.0, 1.0 + 1.0j]]), + pytest.param( + "cos", [-dpnp.pi / 2, -dpnp.pi / 4, 0.0, dpnp.pi / 4, dpnp.pi / 2] + ), + pytest.param("cosh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("copy", [1.0, 2.0, 3.0]), pytest.param("cumprod", [[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]), pytest.param("cumsum", [[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]), @@ -239,11 +249,20 @@ def test_meshgrid(device_x, device_y): pytest.param("nancumsum", [1.0, dpnp.nan]), pytest.param("nanprod", [1.0, dpnp.nan]), pytest.param("nansum", [1.0, dpnp.nan]), - pytest.param("negative", [1.0, -1.0]), + pytest.param("negative", [1.0, 0.0, -1.0]), pytest.param("prod", [1.0, 2.0]), - pytest.param("sign", [-5.0, 4.5]), + pytest.param("sign", [-5.0, 0.0, 4.5]), + pytest.param("signbit", [-5.0, 0.0, 4.5]), + pytest.param( + "sin", [-dpnp.pi / 2, -dpnp.pi / 4, 0.0, dpnp.pi / 4, dpnp.pi / 2] + ), + pytest.param("sinh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("sqrt", [1.0, 3.0, 9.0]), pytest.param("sum", [1.0, 2.0]), + pytest.param( + "tan", [-dpnp.pi / 2, -dpnp.pi / 4, 0.0, dpnp.pi / 4, dpnp.pi / 2] + ), + pytest.param("tanh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("trapz", [[0.0, 1.0, 2.0], [3.0, 4.0, 5.0]]), pytest.param("trunc", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), ], @@ -254,13 +273,14 @@ def test_meshgrid(device_x, device_y): ids=[device.filter_string for device in valid_devices], ) def test_1in_1out(func, data, device): - x_orig = numpy.array(data) - expected = getattr(numpy, func)(x_orig) - x = dpnp.array(data, device=device) result = getattr(dpnp, func)(x) - assert_allclose(result, expected) + x_orig = dpnp.asnumpy(x) + expected = getattr(numpy, func)(x_orig) + + tol = numpy.finfo(x.dtype).resolution + assert_allclose(result, expected, rtol=tol) expected_queue = x.get_array().sycl_queue result_queue = result.get_array().sycl_queue @@ -268,6 +288,35 @@ def test_1in_1out(func, data, device): assert_sycl_queue_equal(result_queue, expected_queue) +@pytest.mark.parametrize( + "device", + valid_devices, + ids=[device.filter_string for device in valid_devices], +) +def test_proj(device): + X = [ + complex(1, 2), + complex(dpnp.inf, -1), + complex(0, -dpnp.inf), + complex(-dpnp.inf, dpnp.nan), + ] + Y = [ + complex(1, 2), + complex(dpnp.inf, -0.0), + complex(dpnp.inf, -0.0), + complex(dpnp.inf, 0.0), + ] + + x = dpnp.array(X, device=device) + result = dpnp.proj(x) + expected = dpnp.array(Y) + assert_allclose(result, expected) + + expected_queue = x.get_array().sycl_queue + result_queue = result.get_array().sycl_queue + assert_sycl_queue_equal(result_queue, expected_queue) + + @pytest.mark.parametrize( "func,data1,data2", [ @@ -276,6 +325,16 @@ def test_1in_1out(func, data, device): [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0], [0.0, 1.0, 2.0, 0.0, 1.0, 2.0, 0.0, 1.0, 2.0], ), + pytest.param( + "allclose", + [1.0, dpnp.inf, -dpnp.inf], + [1.0, dpnp.inf, -dpnp.inf], + ), + pytest.param( + "arctan2", + [[-1, +1, +1, -1]], + [[-1, -1, +1, +1]], + ), pytest.param("copysign", [0.0, 1.0, 2.0], [-1.0, 0.0, 1.0]), pytest.param("cross", [1.0, 2.0, 3.0], [4.0, 5.0, 6.0]), pytest.param( @@ -338,7 +397,7 @@ def test_2in_1out(func, data1, data2, device): x2 = dpnp.array(data2, device=device) result = getattr(dpnp, func)(x1, x2) - assert_array_equal(result, expected) + assert_allclose(result, expected) assert_sycl_queue_equal(result.sycl_queue, x1.sycl_queue) assert_sycl_queue_equal(result.sycl_queue, x2.sycl_queue) @@ -603,7 +662,7 @@ def test_out_2in_1out(func, data1, data2, device): result = dpnp.empty_like(dp_out) getattr(dpnp, func)(x1, x2, out=result) - assert_array_equal(result, expected) + assert_allclose(result, expected) assert_sycl_queue_equal(result.sycl_queue, x1.sycl_queue) assert_sycl_queue_equal(result.sycl_queue, x2.sycl_queue) @@ -731,11 +790,12 @@ def test_eig(device): ) size = 4 - a = numpy.arange(size * size, dtype="float64").reshape((size, size)) + dtype = dpnp.default_float_type(device) + a = numpy.arange(size * size, dtype=dtype).reshape((size, size)) symm_orig = ( numpy.tril(a) + numpy.tril(a, -1).T - + numpy.diag(numpy.full((size,), size * size, dtype="float64")) + + numpy.diag(numpy.full((size,), size * size, dtype=dtype)) ) numpy_data = symm_orig dpnp_symm_orig = dpnp.array(numpy_data, device=device) @@ -780,11 +840,12 @@ def test_eig(device): ) def test_eigh(device): size = 4 - a = numpy.arange(size * size, dtype=numpy.float64).reshape((size, size)) + dtype = dpnp.default_float_type(device) + a = numpy.arange(size * size, dtype=dtype).reshape((size, size)) symm_orig = ( numpy.tril(a) + numpy.tril(a, -1).T - + numpy.diag(numpy.full((size,), size * size, dtype=numpy.float64)) + + numpy.diag(numpy.full((size,), size * size, dtype=dtype)) ) numpy_data = symm_orig dpnp_symm_orig = dpnp.array(numpy_data, device=device) @@ -876,10 +937,9 @@ def test_matrix_rank(device): ids=[device.filter_string for device in valid_devices], ) def test_qr(device): - tol = 1e-11 - data = [[1, 2, 3], [1, 2, 3]] - numpy_data = numpy.array(data) + data = [[1.0, 2.0, 3.0], [1.0, 2.0, 3.0]] dpnp_data = dpnp.array(data, device=device) + numpy_data = numpy.array(data, dtype=dpnp_data.dtype) np_q, np_r = numpy.linalg.qr(numpy_data, "reduced") dpnp_q, dpnp_r = dpnp.linalg.qr(dpnp_data, "reduced") @@ -889,8 +949,8 @@ def test_qr(device): assert dpnp_q.shape == np_q.shape assert dpnp_r.shape == np_r.shape - assert_allclose(dpnp_q, np_q, rtol=tol, atol=tol) - assert_allclose(dpnp_r, np_r, rtol=tol, atol=tol) + assert_dtype_allclose(dpnp_q, np_q) + assert_dtype_allclose(dpnp_r, np_r) expected_queue = dpnp_data.get_array().sycl_queue dpnp_q_queue = dpnp_q.get_array().sycl_queue @@ -907,10 +967,13 @@ def test_qr(device): ids=[device.filter_string for device in valid_devices], ) def test_svd(device): - tol = 1e-12 shape = (2, 2) - numpy_data = numpy.arange(shape[0] * shape[1]).reshape(shape) - dpnp_data = dpnp.arange(shape[0] * shape[1], device=device).reshape(shape) + dtype = dpnp.default_float_type(device) + numpy_data = numpy.arange(shape[0] * shape[1], dtype=dtype).reshape(shape) + dpnp_data = dpnp.arange( + shape[0] * shape[1], dtype=dtype, device=device + ).reshape(shape) + np_u, np_s, np_vt = numpy.linalg.svd(numpy_data) dpnp_u, dpnp_s, dpnp_vt = dpnp.linalg.svd(dpnp_data) @@ -927,11 +990,8 @@ def test_svd(device): dpnp_diag_s[i, i] = dpnp_s[i] # check decomposition - assert_allclose( - dpnp_data, - dpnp.dot(dpnp_u, dpnp.dot(dpnp_diag_s, dpnp_vt)), - rtol=tol, - atol=tol, + assert_dtype_allclose( + dpnp_data, dpnp.dot(dpnp_u, dpnp.dot(dpnp_diag_s, dpnp_vt)) ) for i in range(min(shape[0], shape[1])): @@ -940,13 +1000,9 @@ def test_svd(device): np_vt[i, :] = -np_vt[i, :] # compare vectors for non-zero values - for i in range(numpy.count_nonzero(np_s > tol)): - assert_allclose( - dpnp.asnumpy(dpnp_u)[:, i], np_u[:, i], rtol=tol, atol=tol - ) - assert_allclose( - dpnp.asnumpy(dpnp_vt)[i, :], np_vt[i, :], rtol=tol, atol=tol - ) + for i in range(numpy.count_nonzero(np_s)): + assert_dtype_allclose(dpnp_u[:, i], np_u[:, i]) + assert_dtype_allclose(dpnp_vt[i, :], np_vt[i, :]) expected_queue = dpnp_data.get_array().sycl_queue dpnp_u_queue = dpnp_u.get_array().sycl_queue @@ -972,7 +1028,7 @@ def test_svd(device): def test_to_device(device_from, device_to): data = [1.0, 1.0, 1.0, 1.0, 1.0] - x = dpnp.array(data, device=device_from) + x = dpnp.array(data, dtype=dpnp.float32, device=device_from) y = x.to_device(device_to) assert y.get_array().sycl_device == device_to @@ -985,7 +1041,14 @@ def test_to_device(device_from, device_to): ) @pytest.mark.parametrize( "func", - ["array", "asarray", "asanyarray", "ascontiguousarray", "asfortranarray"], + [ + "array", + "asarray", + "asanyarray", + "ascontiguousarray", + "asfarray", + "asfortranarray", + ], ) @pytest.mark.parametrize( "device_param", ["", "None", "sycl_device"], ids=["Empty", "None", "device"] @@ -1071,3 +1134,25 @@ def test_asarray(device_x, device_y): x = dpnp.array([1, 2, 3], device=device_x) y = dpnp.asarray([x, x, x], device=device_y) assert_sycl_queue_equal(y.sycl_queue, x.to_device(device_y).sycl_queue) + + +@pytest.mark.parametrize( + "device", + valid_devices, + ids=[device.filter_string for device in valid_devices], +) +def test_take(device): + numpy_data = numpy.arange(5) + dpnp_data = dpnp.array(numpy_data, device=device) + + ind = [0, 2, 4] + dpnp_ind = dpnp.array(ind, device=device) + + result = dpnp.take(dpnp_data, dpnp_ind) + expected = numpy.take(numpy_data, ind) + assert_allclose(expected, result) + + expected_queue = dpnp_data.get_array().sycl_queue + result_queue = result.get_array().sycl_queue + + assert_sycl_queue_equal(result_queue, expected_queue) diff --git a/tests/test_umath.py b/tests/test_umath.py index 8f1ac15d6ad8..791e9eb30223 100644 --- a/tests/test_umath.py +++ b/tests/test_umath.py @@ -11,6 +11,7 @@ from .helper import ( get_all_dtypes, get_complex_dtypes, + get_float_dtypes, has_support_aspect16, has_support_aspect64, ) @@ -175,6 +176,89 @@ def test_invalid_shape(self, shape): dpnp.sin(dp_array, out=dp_out) +class TestSinh: + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_complex=True) + ) + def test_sinh(self, dtype): + np_array = numpy.arange(10, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.float64) + + # DPNP + dp_out_dtype = dpnp.float32 + if has_support_aspect64() and dtype != dpnp.float32: + dp_out_dtype = dpnp.float64 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) + result = dpnp.sinh(dp_array, out=dp_out) + + # original + expected = numpy.sinh(np_array, out=np_out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.parametrize("dtype", get_complex_dtypes()) + def test_sinh_complex(self, dtype): + np_array = numpy.arange(10, 20, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.complex128) + + # DPNP + dp_out_dtype = dpnp.complex64 + if has_support_aspect64() and dtype != dpnp.complex64: + dp_out_dtype = dpnp.complex128 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) + result = dpnp.sinh(dp_array, out=dp_out) + + # original + expected = numpy.sinh(np_array, out=np_out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.usefixtures("suppress_divide_numpy_warnings") + @pytest.mark.skipif( + not has_support_aspect16(), reason="No fp16 support by device" + ) + def test_sinh_bool(self): + np_array = numpy.arange(2, dtype=numpy.bool_) + np_out = numpy.empty(2, dtype=numpy.float16) + + # DPNP + dp_array = dpnp.array(np_array, dtype=np_array.dtype) + dp_out = dpnp.array(np_out, dtype=np_out.dtype) + result = dpnp.sinh(dp_array, out=dp_out) + + # original + expected = numpy.sinh(np_array, out=np_out) + assert_allclose(expected, result) + + @pytest.mark.parametrize( + "dtype", + [numpy.float32, numpy.int64, numpy.int32], + ids=["numpy.float32", "numpy.int64", "numpy.int32"], + ) + def test_invalid_dtype(self, dtype): + dp_array = dpnp.arange(10, dtype=dpnp.complex64) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.sinh(dp_array, out=dp_out) + + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape): + dp_array = dpnp.arange(10) + dp_out = dpnp.empty(shape, dtype=dp_array.dtype) + + with pytest.raises(ValueError): + dpnp.sinh(dp_array, out=dp_out) + + class TestCos: @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True) @@ -258,6 +342,89 @@ def test_invalid_shape(self, shape): dpnp.cos(dp_array, out=dp_out) +class TestCosh: + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_complex=True) + ) + def test_cosh(self, dtype): + np_array = numpy.arange(10, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.float64) + + # DPNP + dp_out_dtype = dpnp.float32 + if has_support_aspect64() and dtype != dpnp.float32: + dp_out_dtype = dpnp.float64 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) + result = dpnp.cosh(dp_array, out=dp_out) + + # original + expected = numpy.cosh(np_array, out=np_out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.parametrize("dtype", get_complex_dtypes()) + def test_cosh_complex(self, dtype): + np_array = numpy.arange(10, 20, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.complex128) + + # DPNP + dp_out_dtype = dpnp.complex64 + if has_support_aspect64() and dtype != dpnp.complex64: + dp_out_dtype = dpnp.complex128 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) + result = dpnp.cosh(dp_array, out=dp_out) + + # original + expected = numpy.cosh(np_array, out=np_out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.usefixtures("suppress_divide_numpy_warnings") + @pytest.mark.skipif( + not has_support_aspect16(), reason="No fp16 support by device" + ) + def test_cosh_bool(self): + np_array = numpy.arange(2, dtype=numpy.bool_) + np_out = numpy.empty(2, dtype=numpy.float16) + + # DPNP + dp_array = dpnp.array(np_array, dtype=np_array.dtype) + dp_out = dpnp.array(np_out, dtype=np_out.dtype) + result = dpnp.cosh(dp_array, out=dp_out) + + # original + expected = numpy.cosh(np_array, out=np_out) + assert_allclose(expected, result) + + @pytest.mark.parametrize( + "dtype", + [numpy.float32, numpy.int64, numpy.int32], + ids=["numpy.float32", "numpy.int64", "numpy.int32"], + ) + def test_invalid_dtype(self, dtype): + dp_array = dpnp.arange(10, dtype=dpnp.complex64) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.cosh(dp_array, out=dp_out) + + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape): + dp_array = dpnp.arange(10) + dp_out = dpnp.empty(shape, dtype=dp_array.dtype) + + with pytest.raises(ValueError): + dpnp.cosh(dp_array, out=dp_out) + + class TestsLog: @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True) @@ -380,161 +547,341 @@ def test_invalid_shape(self, shape): dpnp.exp(dp_array, out=dp_out) +class TestArccos: + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") + def test_arccos(self, dtype): + array_data = numpy.arange(-9, 10, 2) / 10 + out = numpy.empty(10, dtype=dtype) + + # DPNP + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) + result = dpnp.arccos(dp_array, out=dp_out) + + # original + np_array = numpy.array(array_data, dtype=dtype) + expected = numpy.arccos(np_array, out=out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] + ) + def test_invalid_dtype(self, dtype): + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.arccos(dp_array, out=dp_out) + + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) + + with pytest.raises(ValueError): + dpnp.arccos(dp_array, out=dp_out) + + +class TestArccosh: + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") + def test_arccosh(self, dtype): + array_data = numpy.arange(2, 12) + out = numpy.empty(10, dtype=dtype) + + # DPNP + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) + result = dpnp.arccosh(dp_array, out=dp_out) + + # original + np_array = numpy.array(array_data, dtype=dtype) + expected = numpy.arccosh(np_array, out=out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] + ) + def test_invalid_dtype(self, dtype): + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.arccosh(dp_array, out=dp_out) + + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) + + with pytest.raises(ValueError): + dpnp.arccosh(dp_array, out=dp_out) + + class TestArcsin: - def test_arcsin(self): - array_data = numpy.arange(10) - out = numpy.empty(10, dtype=numpy.float64) + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") + def test_arcsin(self, dtype): + array_data = numpy.arange(-9, 10, 2) / 10 + out = numpy.empty(10, dtype=dtype) # DPNP - dp_array = dpnp.array(array_data, dtype=dpnp.float64) - dp_out = dpnp.array(out, dtype=dpnp.float64) + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) result = dpnp.arcsin(dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=numpy.float64) + np_array = numpy.array(array_data, dtype=dtype) expected = numpy.arcsin(np_array, out=out) - assert_array_equal(expected, result) + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) @pytest.mark.parametrize( - "dtype", - [numpy.float32, numpy.int64, numpy.int32], - ids=["numpy.float32", "numpy.int64", "numpy.int32"], + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] ) def test_invalid_dtype(self, dtype): - dp_array = dpnp.arange(10, dtype=dpnp.float64) + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) dp_out = dpnp.empty(10, dtype=dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.arcsin(dp_array, out=dp_out) + @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.parametrize( "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) - def test_invalid_shape(self, shape): - dp_array = dpnp.arange(10, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) with pytest.raises(ValueError): dpnp.arcsin(dp_array, out=dp_out) +class TestArcsinh: + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") + def test_arcsinh(self, dtype): + array_data = numpy.arange(10) + out = numpy.empty(10, dtype=dtype) + + # DPNP + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) + result = dpnp.arcsinh(dp_array, out=dp_out) + + # original + np_array = numpy.array(array_data, dtype=dtype) + expected = numpy.arcsinh(np_array, out=out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] + ) + def test_invalid_dtype(self, dtype): + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.arcsinh(dp_array, out=dp_out) + + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) + + with pytest.raises(ValueError): + dpnp.arcsinh(dp_array, out=dp_out) + + class TestArctan: - def test_arctan(self): + @pytest.mark.parametrize("dtype", get_float_dtypes()) + def test_arctan(self, dtype): array_data = numpy.arange(10) - out = numpy.empty(10, dtype=numpy.float64) + out = numpy.empty(10, dtype=dtype) # DPNP - dp_array = dpnp.array(array_data, dtype=dpnp.float64) - dp_out = dpnp.array(out, dtype=dpnp.float64) + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) result = dpnp.arctan(dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=numpy.float64) + np_array = numpy.array(array_data, dtype=dtype) expected = numpy.arctan(np_array, out=out) - assert_array_equal(expected, result) + tol = numpy.finfo(dtype).resolution + assert_allclose(expected, result, tol) @pytest.mark.parametrize( - "dtype", - [numpy.float32, numpy.int64, numpy.int32], - ids=["numpy.float32", "numpy.int64", "numpy.int32"], + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] ) def test_invalid_dtype(self, dtype): - dp_array = dpnp.arange(10, dtype=dpnp.float64) + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) dp_out = dpnp.empty(10, dtype=dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.arctan(dp_array, out=dp_out) + @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.parametrize( "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) - def test_invalid_shape(self, shape): - dp_array = dpnp.arange(10, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) with pytest.raises(ValueError): dpnp.arctan(dp_array, out=dp_out) +class TestArctanh: + @pytest.mark.parametrize("dtype", get_float_dtypes()) + def test_arctanh(self, dtype): + array_data = numpy.arange(-9, 10, 2) / 10 + out = numpy.empty(10, dtype=dtype) + + # DPNP + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) + result = dpnp.arctanh(dp_array, out=dp_out) + + # original + np_array = numpy.array(array_data, dtype=dtype) + expected = numpy.arctanh(np_array, out=out) + + tol = numpy.finfo(dtype).resolution + assert_allclose(expected, result, tol) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] + ) + def test_invalid_dtype(self, dtype): + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.arctanh(dp_array, out=dp_out) + + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) + + with pytest.raises(ValueError): + dpnp.arctanh(dp_array, out=dp_out) + + class TestTan: - def test_tan(self): + @pytest.mark.parametrize("dtype", get_float_dtypes()) + def test_tan(self, dtype): array_data = numpy.arange(10) - out = numpy.empty(10, dtype=numpy.float64) + out = numpy.empty(10, dtype=dtype) # DPNP - dp_array = dpnp.array(array_data, dtype=dpnp.float64) - dp_out = dpnp.array(out, dtype=dpnp.float64) + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) result = dpnp.tan(dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=numpy.float64) + np_array = numpy.array(array_data, dtype=dtype) expected = numpy.tan(np_array, out=out) - assert_allclose(expected, result) + tol = numpy.finfo(dtype).resolution + assert_allclose(expected, result, rtol=tol) @pytest.mark.parametrize( - "dtype", - [numpy.float32, numpy.int64, numpy.int32], - ids=["numpy.float32", "numpy.int64", "numpy.int32"], + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] ) def test_invalid_dtype(self, dtype): - dp_array = dpnp.arange(10, dtype=dpnp.float64) + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) dp_out = dpnp.empty(10, dtype=dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.tan(dp_array, out=dp_out) + @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.parametrize( "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) - def test_invalid_shape(self, shape): - dp_array = dpnp.arange(10, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) with pytest.raises(ValueError): dpnp.tan(dp_array, out=dp_out) class TestArctan2: - def test_arctan2(self): + @pytest.mark.parametrize("dtype", get_float_dtypes()) + def test_arctan2(self, dtype): array_data = numpy.arange(10) - out = numpy.empty(10, dtype=numpy.float64) + out = numpy.empty(10, dtype=dtype) # DPNP - dp_array = dpnp.array(array_data, dtype=dpnp.float64) - dp_out = dpnp.array(out, dtype=dpnp.float64) + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) result = dpnp.arctan2(dp_array, dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=numpy.float64) + np_array = numpy.array(array_data, dtype=dtype) expected = numpy.arctan2(np_array, np_array, out=out) - assert_array_equal(expected, result) + assert_allclose(expected, result) @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True, no_none=True) ) def test_out_dtypes(self, dtype): + if has_support_aspect64() and dtype != numpy.float32: + dtype_out = numpy.float64 + else: + dtype_out = numpy.float32 size = 2 if dtype == dpnp.bool else 10 np_array = numpy.arange(size, dtype=dtype) - np_out = numpy.empty(size, dtype=numpy.complex64) + np_out = numpy.empty(size, dtype=dtype_out) expected = numpy.arctan2(np_array, np_array, out=np_out) dp_array = dpnp.arange(size, dtype=dtype) - dp_out = dpnp.empty(size, dtype=dpnp.complex64) + dp_out = dpnp.empty(size, dtype=dtype_out) result = dpnp.arctan2(dp_array, dp_array, out=dp_out) assert_allclose(expected, result) + @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.parametrize( "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) - def test_invalid_shape(self, shape): - dp_array = dpnp.arange(10, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) with pytest.raises(ValueError): dpnp.arctan2(dp_array, dp_array, out=dp_out) diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index 4bf16022ba90..a935c699f771 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -83,8 +83,35 @@ def test_coerced_usm_types_remainder(usm_type_x, usm_type_y): y = dp.arange(100, usm_type=usm_type_y).reshape(10, 10) y = y.T + 1 + z = 100 % y + z = y % 7 z = x % y + # inplace remainder + z %= y + z %= 5 + + assert x.usm_type == usm_type_x + assert y.usm_type == usm_type_y + assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_y]) + + +@pytest.mark.parametrize("usm_type_x", list_of_usm_types, ids=list_of_usm_types) +@pytest.mark.parametrize("usm_type_y", list_of_usm_types, ids=list_of_usm_types) +def test_coerced_usm_types_floor_divide(usm_type_x, usm_type_y): + x = dp.arange(100, usm_type=usm_type_x).reshape(10, 10) + y = dp.arange(100, usm_type=usm_type_y).reshape(10, 10) + x = x + 1.5 + y = y.T + 0.5 + + z = 3.4 // y + z = y // 2.7 + z = x // y + + # inplace floor_divide + z //= y + z //= 2.5 + assert x.usm_type == usm_type_x assert y.usm_type == usm_type_y assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_y]) @@ -132,7 +159,14 @@ def test_array_creation(func, args, usm_type_x, usm_type_y): @pytest.mark.parametrize( "func", - ["array", "asarray", "asanyarray", "ascontiguousarray", "asfortranarray"], + [ + "array", + "asarray", + "asanyarray", + "ascontiguousarray", + "asfarray", + "asfortranarray", + ], ) @pytest.mark.parametrize("usm_type_x", list_of_usm_types, ids=list_of_usm_types) @pytest.mark.parametrize("usm_type_y", list_of_usm_types, ids=list_of_usm_types) @@ -248,12 +282,32 @@ def test_meshgrid(usm_type_x, usm_type_y): @pytest.mark.parametrize( "func,data", [ + pytest.param("arccos", [-0.5, 0.0, 0.5]), + pytest.param("arccosh", [1.5, 3.5, 5.0]), + pytest.param("arcsin", [-0.5, 0.0, 0.5]), + pytest.param("arcsinh", [-5.0, -3.5, 0.0, 3.5, 5.0]), + pytest.param("arctan", [-1.0, 0.0, 1.0]), + pytest.param("arctanh", [-0.5, 0.0, 0.5]), pytest.param("ceil", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), pytest.param("conjugate", [[1.0 + 1.0j, 0.0], [0.0, 1.0 + 1.0j]]), + pytest.param( + "cos", [-dp.pi / 2, -dp.pi / 4, 0.0, dp.pi / 4, dp.pi / 2] + ), + pytest.param("cosh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("floor", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), - pytest.param("negative", [1.0, -1.0]), - pytest.param("sign", [-5.0, 4.5]), + pytest.param("negative", [1.0, 0.0, -1.0]), + pytest.param("proj", [complex(1.0, 2.0), complex(dp.inf, -1.0)]), + pytest.param("sign", [-5.0, 0.0, 4.5]), + pytest.param("signbit", [-5.0, 0.0, 4.5]), + pytest.param( + "sin", [-dp.pi / 2, -dp.pi / 4, 0.0, dp.pi / 4, dp.pi / 2] + ), + pytest.param("sinh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("sqrt", [1.0, 3.0, 9.0]), + pytest.param( + "tan", [-dp.pi / 2, -dp.pi / 4, 0.0, dp.pi / 4, dp.pi / 2] + ), + pytest.param("tanh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("trunc", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), ], ) @@ -268,6 +322,16 @@ def test_1in_1out(func, data, usm_type): @pytest.mark.parametrize( "func,data1,data2", [ + pytest.param( + "allclose", + [[1.2, -0.0], [-7, 2.34567]], + [[1.2, 0.0], [-7, 2.34567]], + ), + pytest.param( + "arctan2", + [[-1, +1, +1, -1]], + [[-1, -1, +1, +1]], + ), pytest.param( "dot", [[0.0, 1.0, 2.0], [3.0, 4.0, 5.0]], @@ -292,3 +356,17 @@ def test_broadcast_to(usm_type): x = dp.ones(7, usm_type=usm_type) y = dp.broadcast_to(x, (2, 7)) assert x.usm_type == y.usm_type + + +@pytest.mark.parametrize("usm_type_x", list_of_usm_types, ids=list_of_usm_types) +@pytest.mark.parametrize( + "usm_type_ind", list_of_usm_types, ids=list_of_usm_types +) +def test_take(usm_type_x, usm_type_ind): + x = dp.arange(5, usm_type=usm_type_x) + ind = dp.array([0, 2, 4], usm_type=usm_type_ind) + z = dp.take(x, ind) + + assert x.usm_type == usm_type_x + assert ind.usm_type == usm_type_ind + assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_ind]) diff --git a/tests/third_party/cupy/core_tests/test_core_elementwise.py b/tests/third_party/cupy/core_tests/test_core_elementwise.py new file mode 100644 index 000000000000..354a56fc5c42 --- /dev/null +++ b/tests/third_party/cupy/core_tests/test_core_elementwise.py @@ -0,0 +1,143 @@ +import unittest + +import numpy +import pytest + +import dpnp as cupy +from tests.third_party.cupy import testing + + +class TestElementwise(unittest.TestCase): + def check_copy(self, dtype, src_id, dst_id): + with cuda.Device(src_id): + src = testing.shaped_arange((2, 3, 4), dtype=dtype) + with cuda.Device(dst_id): + dst = cupy.empty((2, 3, 4), dtype=dtype) + _core.elementwise_copy(src, dst) + testing.assert_allclose(src, dst) + + @pytest.mark.skip("`device` argument isn't supported") + @testing.for_all_dtypes() + def test_copy(self, dtype): + device_id = cuda.Device().id + self.check_copy(dtype, device_id, device_id) + + @pytest.mark.skip("`device` argument isn't supported") + @testing.for_all_dtypes() + def test_copy_multigpu_nopeer(self, dtype): + if cuda.runtime.deviceCanAccessPeer(0, 1) == 1: + pytest.skip("peer access is available") + with self.assertRaises(ValueError): + self.check_copy(dtype, 0, 1) + + @pytest.mark.skip("`device` argument isn't supported") + @testing.for_all_dtypes() + def test_copy_multigpu_peer(self, dtype): + if cuda.runtime.deviceCanAccessPeer(0, 1) != 1: + pytest.skip("peer access is unavailable") + with pytest.warns(cupy._util.PerformanceWarning): + self.check_copy(dtype, 0, 1) + + @testing.for_orders("CFAK") + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose() + def test_copy_zero_sized_array1(self, xp, dtype, order): + src = xp.empty((0,), dtype=dtype) + res = xp.copy(src, order=order) + assert src is not res + return res + + @testing.for_orders("CFAK") + @testing.for_all_dtypes() + @testing.numpy_cupy_allclose() + def test_copy_zero_sized_array2(self, xp, dtype, order): + src = xp.empty((1, 0, 2), dtype=dtype) + res = xp.copy(src, order=order) + assert src is not res + return res + + @testing.for_orders("CFAK") + def test_copy_orders(self, order): + a = cupy.empty((2, 3, 4)) + b = cupy.copy(a, order) + + a_cpu = numpy.empty((2, 3, 4)) + b_cpu = numpy.copy(a_cpu, order) + + assert b.strides == tuple(x / b_cpu.itemsize for x in b_cpu.strides) + + +@pytest.mark.skip("`ElementwiseKernel` function isn't supported") +class TestElementwiseInvalidShape(unittest.TestCase): + def test_invalid_shape(self): + with self.assertRaisesRegex(ValueError, "Out shape is mismatched"): + f = cupy.ElementwiseKernel("T x", "T y", "y += x") + x = cupy.arange(12).reshape(3, 4) + y = cupy.arange(4) + f(x, y) + + +@pytest.mark.skip("`ElementwiseKernel` function isn't supported") +class TestElementwiseInvalidArgument(unittest.TestCase): + def test_invalid_kernel_name(self): + with self.assertRaisesRegex(ValueError, "Invalid kernel name"): + cupy.ElementwiseKernel("T x", "", "", "1") + + +@pytest.mark.skip("`iinfo` function isn't supported") +class TestElementwiseType(unittest.TestCase): + @testing.for_int_dtypes(no_bool=True) + @testing.numpy_cupy_array_equal() + def test_large_int_upper_1(self, xp, dtype): + a = xp.array([0], dtype=xp.int8) + b = xp.iinfo(dtype).max + return a + b + + @testing.for_int_dtypes(no_bool=True) + @testing.numpy_cupy_array_equal() + def test_large_int_upper_2(self, xp, dtype): + a = xp.array([1], dtype=xp.int8) + b = xp.iinfo(dtype).max - 1 + return a + b + + @testing.for_int_dtypes(no_bool=True) + @testing.numpy_cupy_array_equal() + def test_large_int_upper_3(self, xp, dtype): + a = xp.array([xp.iinfo(dtype).max], dtype=dtype) + b = xp.int8(0) + return a + b + + @testing.for_int_dtypes(no_bool=True) + @testing.numpy_cupy_array_equal() + def test_large_int_upper_4(self, xp, dtype): + a = xp.array([xp.iinfo(dtype).max - 1], dtype=dtype) + b = xp.int8(1) + return a + b + + @testing.for_int_dtypes(no_bool=True) + @testing.numpy_cupy_array_equal() + def test_large_int_lower_1(self, xp, dtype): + a = xp.array([0], dtype=xp.int8) + b = xp.iinfo(dtype).min + return a + b + + @testing.for_int_dtypes(no_bool=True) + @testing.numpy_cupy_array_equal() + def test_large_int_lower_2(self, xp, dtype): + a = xp.array([-1], dtype=xp.int8) + b = xp.iinfo(dtype).min + 1 + return a + b + + @testing.for_int_dtypes(no_bool=True) + @testing.numpy_cupy_array_equal() + def test_large_int_lower_3(self, xp, dtype): + a = xp.array([xp.iinfo(dtype).min], dtype=dtype) + b = xp.int8(0) + return a + b + + @testing.for_int_dtypes(no_bool=True) + @testing.numpy_cupy_array_equal() + def test_large_int_lower_4(self, xp, dtype): + a = xp.array([xp.iinfo(dtype).min + 1], dtype=dtype) + b = xp.int8(-1) + return a + b diff --git a/tests/third_party/cupy/core_tests/test_ndarray_math.py b/tests/third_party/cupy/core_tests/test_ndarray_math.py index 9b16b76e4c82..5ed98f374bf9 100644 --- a/tests/third_party/cupy/core_tests/test_ndarray_math.py +++ b/tests/third_party/cupy/core_tests/test_ndarray_math.py @@ -4,6 +4,7 @@ import pytest import dpnp as cupy +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing @@ -34,7 +35,8 @@ def test_round(self, xp, dtype): @testing.numpy_cupy_array_equal() def test_round_out(self, xp): - a = testing.shaped_random(self.shape, xp, scale=100, dtype="d") + dtype = "d" if has_support_aspect64() else "f" + a = testing.shaped_random(self.shape, xp, scale=100, dtype=dtype) out = xp.empty_like(a) a.round(self.decimals, out) return out @@ -50,6 +52,7 @@ def test_round_out(self, xp): } ) ) +@pytest.mark.usefixtures("allow_fall_back_on_numpy") class TestRoundHalfway(unittest.TestCase): shape = (20,) diff --git a/tests/third_party/cupy/creation_tests/test_from_data.py b/tests/third_party/cupy/creation_tests/test_from_data.py index be606d1840f3..a517a2cb2a10 100644 --- a/tests/third_party/cupy/creation_tests/test_from_data.py +++ b/tests/third_party/cupy/creation_tests/test_from_data.py @@ -501,7 +501,9 @@ def test_asarray_from_big_endian(self, xp, dtype): # happens to work before the change in #5828 return b + b - @pytest.mark.usefixtures("allow_fall_back_on_numpy") + @pytest.mark.skip( + "TODO: remove once dpctl gh-1376 is merged to gold branch" + ) @testing.for_CF_orders() @testing.for_all_dtypes() @testing.numpy_cupy_array_equal() diff --git a/tests/third_party/cupy/creation_tests/test_ranges.py b/tests/third_party/cupy/creation_tests/test_ranges.py index b95d254c9980..c165b4e94a29 100644 --- a/tests/third_party/cupy/creation_tests/test_ranges.py +++ b/tests/third_party/cupy/creation_tests/test_ranges.py @@ -6,6 +6,7 @@ import pytest import dpnp as cupy +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing @@ -60,7 +61,7 @@ def test_arange9(self): def test_arange_no_dtype_int(self, xp): return xp.arange(1, 11, 2) - @testing.numpy_cupy_array_equal() + @testing.numpy_cupy_allclose(rtol=1e-4, type_check=has_support_aspect64()) def test_arange_no_dtype_float(self, xp): return xp.arange(1.0, 11.0, 2.0) @@ -120,11 +121,11 @@ def test_linspace_with_retstep(self, xp, dtype): self.assertEqual(step, 2.5) return x - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(rtol=1e-4, type_check=has_support_aspect64()) def test_linspace_no_dtype_int(self, xp): return xp.linspace(0, 10, 50) - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(rtol=1e-4, type_check=has_support_aspect64()) def test_linspace_no_dtype_float(self, xp): return xp.linspace(0.0, 10.0, 50) @@ -139,21 +140,23 @@ def test_linspace_neg_num(self): @testing.numpy_cupy_allclose() def test_linspace_float_overflow(self, xp): - return xp.linspace(0.0, sys.float_info.max / 5, 10, dtype=float) + dtype = cupy.default_float_type() + return xp.linspace(0.0, numpy.finfo(dtype).max / 5, 10, dtype=dtype) - @testing.numpy_cupy_array_equal() + @testing.numpy_cupy_allclose() def test_linspace_float_underflow(self, xp): # find minimum subnormal number - x = sys.float_info.min + dtype = cupy.default_float_type() + x = numpy.finfo(dtype).min while x / 2 > 0: x /= 2 - return xp.linspace(0.0, x, 10, dtype=float) + return xp.linspace(0.0, x, 10, dtype=dtype) @testing.with_requires("numpy>=1.16") @testing.for_all_dtypes_combination( names=("dtype_range", "dtype_out"), no_bool=True, no_complex=True ) - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(rtol=1e-04) def test_linspace_array_start_stop(self, xp, dtype_range, dtype_out): start = xp.array([0, 120], dtype=dtype_range) stop = xp.array([100, 0], dtype=dtype_range) @@ -163,7 +166,7 @@ def test_linspace_array_start_stop(self, xp, dtype_range, dtype_out): @testing.for_all_dtypes_combination( names=("dtype_range", "dtype_out"), no_bool=True, no_complex=True ) - @testing.numpy_cupy_array_equal() + @testing.numpy_cupy_allclose(rtol=1e-04) def test_linspace_mixed_start_stop(self, xp, dtype_range, dtype_out): start = 0.0 if xp.dtype(dtype_range).kind in "u": @@ -176,7 +179,7 @@ def test_linspace_mixed_start_stop(self, xp, dtype_range, dtype_out): @testing.for_all_dtypes_combination( names=("dtype_range", "dtype_out"), no_bool=True, no_complex=True ) - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(rtol=1e-04) def test_linspace_mixed_start_stop2(self, xp, dtype_range, dtype_out): if xp.dtype(dtype_range).kind in "u": start = xp.array([160, 120], dtype=dtype_range) @@ -205,7 +208,7 @@ def test_linspace_complex_start_stop(self, xp, dtype): @testing.with_requires("numpy>=1.16") @testing.for_all_dtypes(no_bool=True) - @testing.numpy_cupy_array_equal() + @testing.numpy_cupy_allclose(rtol=1e-04) def test_linspace_start_stop_list(self, xp, dtype): start = [0, 0] stop = [100, 16] @@ -241,12 +244,12 @@ def test_logspace_no_endpoint(self, xp, dtype): return xp.logspace(0, 2, 5, dtype=dtype, endpoint=False) @pytest.mark.usefixtures("allow_fall_back_on_numpy") - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(rtol=1e-4, type_check=has_support_aspect64()) def test_logspace_no_dtype_int(self, xp): return xp.logspace(0, 2) @pytest.mark.usefixtures("allow_fall_back_on_numpy") - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(rtol=1e-4, type_check=has_support_aspect64()) def test_logspace_no_dtype_float(self, xp): return xp.logspace(0.0, 2.0) @@ -262,7 +265,7 @@ def test_logspace_neg_num(self): @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes(no_bool=True) - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(rtol=1e-04) def test_logspace_base(self, xp, dtype): return xp.logspace(0, 2, 5, base=2.0, dtype=dtype) @@ -323,7 +326,7 @@ def test_mgrid0(self, xp): def test_mgrid1(self, xp): return xp.mgrid[-10:10] - @testing.numpy_cupy_array_equal() + @testing.numpy_cupy_allclose(rtol=1e-4, type_check=has_support_aspect64()) def test_mgrid2(self, xp): return xp.mgrid[-10:10:10j] @@ -333,7 +336,7 @@ def test_mgrid3(self, xp): y = xp.ones(10)[:, None] return xp.mgrid[x:y:10j] - @testing.numpy_cupy_array_equal() + @testing.numpy_cupy_allclose(rtol=1e-4, type_check=has_support_aspect64()) def test_mgrid4(self, xp): # check len(keys) > 1 return xp.mgrid[-10:10:10j, -10:10:10j] @@ -356,7 +359,7 @@ def test_ogrid0(self, xp): def test_ogrid1(self, xp): return xp.ogrid[-10:10] - @testing.numpy_cupy_array_equal() + @testing.numpy_cupy_allclose(rtol=1e-4, type_check=has_support_aspect64()) def test_ogrid2(self, xp): return xp.ogrid[-10:10:10j] diff --git a/tests/third_party/cupy/fft_tests/test_fft.py b/tests/third_party/cupy/fft_tests/test_fft.py index 4650fda557ab..98c4ad29288a 100644 --- a/tests/third_party/cupy/fft_tests/test_fft.py +++ b/tests/third_party/cupy/fft_tests/test_fft.py @@ -85,6 +85,7 @@ class TestFft2(unittest.TestCase): atol=1e-7, accept_error=ValueError, type_check=has_support_aspect64(), + contiguous_check=False, ) def test_fft2(self, xp, dtype): a = testing.shaped_random(self.shape, xp, dtype) @@ -98,6 +99,7 @@ def test_fft2(self, xp, dtype): atol=1e-7, accept_error=ValueError, type_check=has_support_aspect64(), + contiguous_check=False, ) def test_ifft2(self, xp, dtype): a = testing.shaped_random(self.shape, xp, dtype) @@ -141,6 +143,7 @@ class TestFftn(unittest.TestCase): atol=1e-7, accept_error=ValueError, type_check=has_support_aspect64(), + contiguous_check=False, ) def test_fftn(self, xp, dtype): a = testing.shaped_random(self.shape, xp, dtype) @@ -154,6 +157,7 @@ def test_fftn(self, xp, dtype): atol=1e-7, accept_error=ValueError, type_check=has_support_aspect64(), + contiguous_check=False, ) def test_ifftn(self, xp, dtype): a = testing.shaped_random(self.shape, xp, dtype) diff --git a/tests/third_party/cupy/indexing_tests/test_insert.py b/tests/third_party/cupy/indexing_tests/test_insert.py index f48cff027335..1840b22d5c46 100644 --- a/tests/third_party/cupy/indexing_tests/test_insert.py +++ b/tests/third_party/cupy/indexing_tests/test_insert.py @@ -71,7 +71,8 @@ def test_place_shape_unmatch_error(self, dtype): { "shape": [(7,), (2, 3), (4, 3, 2)], "mode": ["raise", "wrap", "clip"], - "n_vals": [0, 1, 3, 4, 5], + # The vals shape of array must be broadcastable to the shape of indices + "n_vals": [1, 4], } ) ) @@ -83,12 +84,17 @@ class TestPut(unittest.TestCase): def test_put(self, xp, dtype): a = testing.shaped_arange(self.shape, xp, dtype) # Take care so that actual indices don't overlap. - if self.mode == "raise": + if self.mode in ("raise"): inds = xp.array([2, -1, 3, 0]) + # `wrap` mode in dpctl.tensor.put is different from numpy.put (#1365): + # numpy`s `wrap` mode wraps indices around for cyclic operations + # while dpctl`s `wrap` mode restricts indices to stay within the array bounds (-n <= i < n). + elif self.mode in ("wrap"): + inds = xp.array([2, -1, 3, -6]) else: inds = xp.array([2, -8, 3, 7]) vals = testing.shaped_random((self.n_vals,), xp, dtype) - xp.put(a, inds, vals, self.mode) + xp.put(a, inds, vals, mode=self.mode) return a @@ -99,9 +105,9 @@ def test_put(self, xp, dtype): } ) ) -@pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.gpu class TestPutScalars(unittest.TestCase): + @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.numpy_cupy_array_equal() def test_put_index_scalar(self, xp): dtype = cupy.float32 diff --git a/tests/third_party/cupy/linalg_tests/test_eigenvalue.py b/tests/third_party/cupy/linalg_tests/test_eigenvalue.py index b5a73c25e95b..99dcfb2127c8 100644 --- a/tests/third_party/cupy/linalg_tests/test_eigenvalue.py +++ b/tests/third_party/cupy/linalg_tests/test_eigenvalue.py @@ -4,22 +4,19 @@ import pytest import dpnp as cupy +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing def _get_hermitian(xp, a, UPLO): - # TODO: remove wrapping, but now there is no dpnp_array.swapaxes() - a = _wrap_as_numpy_array(xp, a) - _xp = numpy - if UPLO == "U": - _a = _xp.triu(a) + _xp.triu(a, k=1).swapaxes(-2, -1).conj() + return xp.triu(a) + xp.triu(a, k=1).swapaxes(-2, -1).conj() else: - _a = _xp.tril(a) + _xp.tril(a, k=-1).swapaxes(-2, -1).conj() - return xp.array(_a) + return xp.tril(a) + xp.tril(a, k=-1).swapaxes(-2, -1).conj() -# TODO: remove once all required functionality is supported +# TODO: +# remove once dpnp.dot and dpnp.matmul support complex types def _wrap_as_numpy_array(xp, a): return a.asnumpy() if xp is cupy else a @@ -33,7 +30,12 @@ def _wrap_as_numpy_array(xp, a): ) class TestEigenvalue(unittest.TestCase): @testing.for_all_dtypes() - @testing.numpy_cupy_allclose(rtol=1e-3, atol=1e-4, contiguous_check=False) + @testing.numpy_cupy_allclose( + rtol=1e-3, + atol=1e-4, + type_check=has_support_aspect64(), + contiguous_check=False, + ) def test_eigh(self, xp, dtype): if xp == numpy and dtype == numpy.float16: # NumPy's eigh does not support float16 @@ -190,7 +192,7 @@ def test_eigvalsh_complex_batched(self, xp, dtype): ) class TestEigenvalueEmpty(unittest.TestCase): @testing.for_dtypes("ifdFD") - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_eigh(self, xp, dtype): a = xp.empty(self.shape, dtype=dtype) assert a.size == 0 diff --git a/tests/third_party/cupy/linalg_tests/test_einsum.py b/tests/third_party/cupy/linalg_tests/test_einsum.py index ba8baa58dcf8..613102c49c3f 100644 --- a/tests/third_party/cupy/linalg_tests/test_einsum.py +++ b/tests/third_party/cupy/linalg_tests/test_einsum.py @@ -5,6 +5,7 @@ import pytest import dpnp as cupy +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing @@ -310,7 +311,9 @@ def setUp(self): self.operands = operands @pytest.mark.usefixtures("allow_fall_back_on_numpy") - @testing.numpy_cupy_allclose(contiguous_check=False) + @testing.numpy_cupy_allclose( + type_check=has_support_aspect64(), contiguous_check=False + ) def test_einsum(self, xp): # TODO(kataoka): support memory efficient cupy.einsum with warnings.catch_warnings(record=True) as ws: diff --git a/tests/third_party/cupy/linalg_tests/test_product.py b/tests/third_party/cupy/linalg_tests/test_product.py index 974ea0af3773..93b13c93e873 100644 --- a/tests/third_party/cupy/linalg_tests/test_product.py +++ b/tests/third_party/cupy/linalg_tests/test_product.py @@ -4,6 +4,7 @@ import pytest import dpnp as cupy +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing @@ -101,7 +102,9 @@ def test_dot_with_out(self, xp, dtype_a, dtype_b, dtype_c): class TestCrossProduct(unittest.TestCase): @testing.for_all_dtypes_combination(["dtype_a", "dtype_b"]) # TODO: remove 'contiguous_check=False' once fixed in dpnp.cross() - @testing.numpy_cupy_allclose(contiguous_check=False) + @testing.numpy_cupy_allclose( + type_check=has_support_aspect64(), contiguous_check=False + ) def test_cross(self, xp, dtype_a, dtype_b): if dtype_a == dtype_b == numpy.bool_: # cross does not support bool-bool inputs. diff --git a/tests/third_party/cupy/logic_tests/test_comparison.py b/tests/third_party/cupy/logic_tests/test_comparison.py index a37479c4fe21..739ff9f7c959 100644 --- a/tests/third_party/cupy/logic_tests/test_comparison.py +++ b/tests/third_party/cupy/logic_tests/test_comparison.py @@ -121,15 +121,15 @@ class TestAllclose(unittest.TestCase): @testing.for_all_dtypes() @testing.numpy_cupy_equal() def test_allclose_finite(self, xp, dtype): - a = xp.array([0.9e-5, 1.1e-5, 1000 + 1e-4, 1000 - 1e-4], dtype=dtype) - b = xp.array([0, 0, 1000, 1000], dtype=dtype) + a = xp.array([0.9e-5, 1.1e-5, 1000 + 1e-4, 1000 - 1e-4]).astype(dtype) + b = xp.array([0, 0, 1000, 1000]).astype(dtype) return xp.allclose(a, b) @testing.for_all_dtypes() @testing.numpy_cupy_equal() def test_allclose_min_int(self, xp, dtype): - a = xp.array([0], dtype=dtype) - b = xp.array([numpy.iinfo("i").min], dtype=dtype) + a = xp.array([0]).astype(dtype) + b = xp.array([numpy.iinfo("i").min]).astype(dtype) return xp.allclose(a, b) @testing.for_float_dtypes() @@ -138,24 +138,25 @@ def test_allclose_infinite(self, xp, dtype): nan = float("nan") inf = float("inf") ninf = float("-inf") - a = xp.array([0, nan, nan, 0, inf, ninf], dtype=dtype) - b = xp.array([0, nan, 0, nan, inf, ninf], dtype=dtype) + a = xp.array([0, nan, nan, 0, inf, ninf]).astype(dtype) + b = xp.array([0, nan, 0, nan, inf, ninf]).astype(dtype) return xp.allclose(a, b) @testing.for_float_dtypes() @testing.numpy_cupy_equal() + @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_allclose_infinite_equal_nan(self, xp, dtype): nan = float("nan") inf = float("inf") ninf = float("-inf") - a = xp.array([0, nan, inf, ninf], dtype=dtype) - b = xp.array([0, nan, inf, ninf], dtype=dtype) + a = xp.array([0, nan, inf, ninf]).astype(dtype) + b = xp.array([0, nan, inf, ninf]).astype(dtype) return xp.allclose(a, b, equal_nan=True) @testing.for_all_dtypes() @testing.numpy_cupy_equal() def test_allclose_array_scalar(self, xp, dtype): - a = xp.array([0.9e-5, 1.1e-5], dtype=dtype) + a = xp.array([0.9e-5, 1.1e-5]).astype(dtype) b = xp.dtype(xp.dtype).type(0) return xp.allclose(a, b) diff --git a/tests/third_party/cupy/manipulation_tests/test_basic.py b/tests/third_party/cupy/manipulation_tests/test_basic.py index 0712f88a736e..46c3533a74ff 100644 --- a/tests/third_party/cupy/manipulation_tests/test_basic.py +++ b/tests/third_party/cupy/manipulation_tests/test_basic.py @@ -120,10 +120,7 @@ def test_copyto_where_multidevice_raises(self, dtype): @testing.for_all_dtypes() def test_copyto_noncontinguous(self, dtype): src = testing.shaped_arange((2, 3, 4), cupy, dtype) - # TODO: replace with - # src = src.swapaxes(0, 1) - # once dpnp.ndarray.swapaxes() is fully implemented - src = cupy.swapaxes(src, 0, 1) + src = src.swapaxes(0, 1) dst = cupy.empty_like(src) cupy.copyto(dst, src) diff --git a/tests/third_party/cupy/manipulation_tests/test_dims.py b/tests/third_party/cupy/manipulation_tests/test_dims.py index 6d59f0640512..461c0bb2ec2b 100644 --- a/tests/third_party/cupy/manipulation_tests/test_dims.py +++ b/tests/third_party/cupy/manipulation_tests/test_dims.py @@ -18,12 +18,10 @@ def check_atleast(self, func, xp): f = numpy.float32(1) return func(a, b, c, d, e, f) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.numpy_cupy_array_equal() def test_atleast_1d1(self, xp): return self.check_atleast(xp.atleast_1d, xp) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.numpy_cupy_array_equal() def test_atleast_1d2(self, xp): a = testing.shaped_arange((1, 3, 2), xp) diff --git a/tests/third_party/cupy/manipulation_tests/test_join.py b/tests/third_party/cupy/manipulation_tests/test_join.py index 054ada455aee..8bffce98d875 100644 --- a/tests/third_party/cupy/manipulation_tests/test_join.py +++ b/tests/third_party/cupy/manipulation_tests/test_join.py @@ -270,14 +270,12 @@ def test_dstack_single_element_3(self, xp): a = testing.shaped_arange((1,), xp) return xp.dstack((a,)) - @pytest.mark.skip("dpnp.hstack() is not implemented yet") @testing.numpy_cupy_array_equal() def test_hstack_vectors(self, xp): a = xp.arange(3) b = xp.arange(2, -1, -1) return xp.hstack((a, b)) - @pytest.mark.skip("dpnp.hstack() is not implemented yet") @testing.numpy_cupy_array_equal() def test_hstack_scalars(self, xp): a = testing.shaped_arange((), xp) @@ -285,7 +283,6 @@ def test_hstack_scalars(self, xp): c = testing.shaped_arange((), xp) return xp.hstack((a, b, c)) - @pytest.mark.skip("dpnp.hstack() is not implemented yet") @testing.numpy_cupy_array_equal() def test_hstack(self, xp): a = testing.shaped_arange((2, 1), xp) @@ -293,7 +290,7 @@ def test_hstack(self, xp): c = testing.shaped_arange((2, 3), xp) return xp.hstack((a, b, c)) - @pytest.mark.skip("dpnp.hstack() is not implemented yet") + @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.with_requires("numpy>=1.24.0") @testing.for_all_dtypes_combination(names=["dtype1", "dtype2"]) @testing.numpy_cupy_array_equal(accept_error=TypeError) @@ -302,18 +299,9 @@ def test_hstack_dtype(self, xp, dtype1, dtype2): b = testing.shaped_arange((3, 4), xp, dtype1) return xp.hstack((a, b), dtype=dtype2) - @pytest.mark.skip("dpnp.hstack() is not implemented yet") + @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.with_requires("numpy>=1.24.0") - @pytest.mark.parametrize( - "casting", - [ - "no", - "equiv", - "safe", - "same_kind", - "unsafe", - ], - ) + @testing.for_castings() @testing.for_all_dtypes_combination(names=["dtype1", "dtype2"]) @testing.numpy_cupy_array_equal( accept_error=(TypeError, numpy.ComplexWarning) diff --git a/tests/third_party/cupy/manipulation_tests/test_kind.py b/tests/third_party/cupy/manipulation_tests/test_kind.py new file mode 100644 index 000000000000..1812d77c0af5 --- /dev/null +++ b/tests/third_party/cupy/manipulation_tests/test_kind.py @@ -0,0 +1,143 @@ +import unittest + +import numpy +import pytest + +import dpnp as cupy +from tests.helper import has_support_aspect64 +from tests.third_party.cupy import testing + + +class TestKind(unittest.TestCase): + @pytest.mark.skip("dpnp.asarray_chkfinite() is not implemented yet") + @testing.for_orders("CFAK") + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_asarray_chkfinite(self, xp, dtype, order): + a = [0, 4, 0, 5] + return xp.asarray_chkfinite(a, dtype=dtype, order=order) + + @pytest.mark.skip("dpnp.asarray_chkfinite() is not implemented yet") + @testing.for_orders("CFAK") + @testing.for_all_dtypes(no_bool=True) + def test_asarray_chkfinite_non_finite_vals(self, dtype, order): + a = [-numpy.inf, 0.0, numpy.inf, numpy.nan] + for xp in (numpy, cupy): + if xp.issubdtype(dtype, xp.integer): + error = OverflowError + else: + error = ValueError + with pytest.raises(error): + xp.asarray_chkfinite(a, dtype=dtype, order=order) + + @testing.for_all_dtypes() + def test_asfarray(self, dtype): + a = cupy.asarray([1, 2, 3]) + a_gpu = cupy.asfarray(a, dtype) + a_cpu = numpy.asfarray(a, dtype) + if ( + has_support_aspect64() + or cupy.issubdtype(dtype, cupy.complexfloating) + or cupy.issubdtype(dtype, cupy.floating) + ): + assert a_cpu.dtype == a_gpu.dtype + else: + assert a_cpu.dtype == cupy.float64 + assert a_gpu.dtype == cupy.float32 + + @testing.for_all_dtypes() + def test_asfortranarray1(self, dtype): + def func(xp): + x = xp.zeros((2, 3), dtype=dtype) + ret = xp.asfortranarray(x) + assert x.flags.c_contiguous + assert ret.flags.f_contiguous + + assert func(numpy) == func(cupy) + + @testing.for_all_dtypes() + def test_asfortranarray2(self, dtype): + def func(xp): + x = xp.zeros((2, 3, 4), dtype=dtype) + ret = xp.asfortranarray(x) + assert x.flags.c_contiguous + assert ret.flags.f_contiguous + + assert func(numpy) == func(cupy) + + @testing.for_all_dtypes() + def test_asfortranarray3(self, dtype): + def func(xp): + x = xp.zeros((2, 3, 4), dtype=dtype) + ret = xp.asfortranarray(xp.asfortranarray(x)) + assert x.flags.c_contiguous + assert ret.flags.f_contiguous + + assert func(numpy) == func(cupy) + + @testing.for_all_dtypes() + def test_asfortranarray4(self, dtype): + def func(xp): + x = xp.zeros((2, 3), dtype=dtype) + x = xp.transpose(x, (1, 0)) + ret = xp.asfortranarray(x) + assert ret.flags.f_contiguous + + assert func(numpy) == func(cupy) + + @testing.for_all_dtypes() + def test_asfortranarray5(self, dtype): + def func(xp): + x = testing.shaped_arange((2, 3), xp, dtype) + ret = xp.asfortranarray(x) + assert x.flags.c_contiguous + assert ret.flags.f_contiguous + + assert func(numpy) == func(cupy) + + @pytest.mark.skip("dpnp.require() is not implemented yet") + @testing.for_all_dtypes() + def test_require_flag_check(self, dtype): + possible_flags = [["C_CONTIGUOUS"], ["F_CONTIGUOUS"]] + x = cupy.zeros((2, 3, 4), dtype=dtype) + for flags in possible_flags: + arr = cupy.require(x, dtype, flags) + for parameter in flags: + assert arr.flags[parameter] + assert arr.dtype == dtype + + @pytest.mark.skip("dpnp.require() is not implemented yet") + @testing.for_all_dtypes() + def test_require_owndata(self, dtype): + x = cupy.zeros((2, 3, 4), dtype=dtype) + arr = x.view() + arr = cupy.require(arr, dtype, ["O"]) + assert arr.flags["OWNDATA"] + + @pytest.mark.skip("dpnp.require() is not implemented yet") + @testing.for_all_dtypes() + def test_require_C_and_F_flags(self, dtype): + x = cupy.zeros((2, 3, 4), dtype=dtype) + with pytest.raises(ValueError): + cupy.require(x, dtype, ["C", "F"]) + + @pytest.mark.skip("dpnp.require() is not implemented yet") + @testing.for_all_dtypes() + def test_require_incorrect_requirments(self, dtype): + x = cupy.zeros((2, 3, 4), dtype=dtype) + with pytest.raises(ValueError): + cupy.require(x, dtype, ["W"]) + + @pytest.mark.skip("dpnp.require() is not implemented yet") + @testing.for_all_dtypes() + def test_require_incorrect_dtype(self, dtype): + x = cupy.zeros((2, 3, 4), dtype=dtype) + with pytest.raises(ValueError): + cupy.require(x, "random", "C") + + @pytest.mark.skip("dpnp.require() is not implemented yet") + @testing.for_all_dtypes() + def test_require_empty_requirements(self, dtype): + x = cupy.zeros((2, 3, 4), dtype=dtype) + x = cupy.require(x, dtype, []) + assert x.flags["C_CONTIGUOUS"] diff --git a/tests/third_party/cupy/manipulation_tests/test_rearrange.py b/tests/third_party/cupy/manipulation_tests/test_rearrange.py new file mode 100644 index 000000000000..f97f611ff252 --- /dev/null +++ b/tests/third_party/cupy/manipulation_tests/test_rearrange.py @@ -0,0 +1,263 @@ +import unittest + +import numpy +import pytest + +import dpnp as cupy +from tests.third_party.cupy import testing + + +@testing.parameterize( + {"shape": (10,), "shift": 2, "axis": None}, + {"shape": (5, 2), "shift": 1, "axis": None}, + {"shape": (5, 2), "shift": -2, "axis": None}, + {"shape": (5, 2), "shift": 1, "axis": 0}, + {"shape": (5, 2), "shift": 1, "axis": -1}, + {"shape": (10,), "shift": 35, "axis": None}, + {"shape": (5, 2), "shift": 11, "axis": 0}, + {"shape": (), "shift": 5, "axis": None}, + {"shape": (5, 2), "shift": 1, "axis": (0, 1)}, + {"shape": (5, 2), "shift": 1, "axis": (0, 0)}, + {"shape": (5, 2), "shift": 50, "axis": 0}, + {"shape": (5, 2), "shift": (2, 1), "axis": (0, 1)}, + {"shape": (5, 2), "shift": (2, 1), "axis": (0, -1)}, + {"shape": (5, 2), "shift": (2, 1), "axis": (1, -1)}, + {"shape": (5, 2), "shift": (2, 1, 3), "axis": 0}, + {"shape": (5, 2), "shift": (2, 1, 3), "axis": None}, +) +class TestRoll(unittest.TestCase): + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_roll(self, xp, dtype): + x = testing.shaped_arange(self.shape, xp, dtype) + return xp.roll(x, self.shift, axis=self.axis) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_roll_cupy_shift(self, xp, dtype): + x = testing.shaped_arange(self.shape, xp, dtype) + shift = self.shift + return xp.roll(x, shift, axis=self.axis) + + +class TestRollTypeError(unittest.TestCase): + def test_roll_invalid_shift(self): + for xp in (numpy, cupy): + x = testing.shaped_arange((5, 2), xp) + with pytest.raises(TypeError): + xp.roll(x, "0", axis=0) + + def test_roll_invalid_axis_type(self): + for xp in (numpy, cupy): + x = testing.shaped_arange((5, 2), xp) + with pytest.raises(TypeError): + xp.roll(x, 2, axis="0") + + +@testing.parameterize( + {"shape": (5, 2, 3), "shift": (2, 2, 2), "axis": (0, 1)}, + {"shape": (5, 2), "shift": 1, "axis": 2}, + {"shape": (5, 2), "shift": 1, "axis": -3}, + {"shape": (5, 2, 2), "shift": (1, 0), "axis": (0, 1, 2)}, + {"shape": (5, 2), "shift": 1, "axis": -3}, + {"shape": (5, 2), "shift": 1, "axis": (1, -3)}, +) +class TestRollValueError(unittest.TestCase): + def test_roll_invalid(self): + for xp in (numpy, cupy): + x = testing.shaped_arange(self.shape, xp) + with pytest.raises(ValueError): + xp.roll(x, self.shift, axis=self.axis) + + def test_roll_invalid_cupy_shift(self): + for xp in (numpy, cupy): + x = testing.shaped_arange(self.shape, xp) + shift = self.shift + with pytest.raises(ValueError): + xp.roll(x, shift, axis=self.axis) + + +class TestFliplr(unittest.TestCase): + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_fliplr_2(self, xp, dtype): + x = testing.shaped_arange((3, 4), xp, dtype) + return xp.fliplr(x) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_fliplr_3(self, xp, dtype): + x = testing.shaped_arange((3, 4, 2), xp, dtype) + return xp.fliplr(x) + + @testing.for_all_dtypes() + def test_fliplr_insufficient_ndim(self, dtype): + for xp in (numpy, cupy): + x = testing.shaped_arange((3,), xp, dtype) + with pytest.raises(ValueError): + xp.fliplr(x) + + +class TestFlipud(unittest.TestCase): + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flipud_1(self, xp, dtype): + x = testing.shaped_arange((3,), xp, dtype) + return xp.flipud(x) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flipud_2(self, xp, dtype): + x = testing.shaped_arange((3, 4), xp, dtype) + return xp.flipud(x) + + @testing.for_all_dtypes() + def test_flipud_insufficient_ndim(self, dtype): + for xp in (numpy, cupy): + x = testing.shaped_arange((), xp, dtype) + with pytest.raises(ValueError): + xp.flipud(x) + + +class TestFlip(unittest.TestCase): + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flip_1(self, xp, dtype): + x = testing.shaped_arange((3,), xp, dtype) + return xp.flip(x, 0) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flip_2(self, xp, dtype): + x = testing.shaped_arange((3, 4), xp, dtype) + return xp.flip(x, 1) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flip_all(self, xp, dtype): + x = testing.shaped_arange((3, 4), xp, dtype) + return xp.flip(x) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flip_with_negative_axis(self, xp, dtype): + x = testing.shaped_arange((3, 4, 2), xp, dtype) + return xp.flip(x, -1) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flip_with_axis_tuple(self, xp, dtype): + x = testing.shaped_arange((3, 4, 2), xp, dtype) + return xp.flip(x, (0, 2)) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flip_dim_0(self, xp, dtype): + x = testing.shaped_arange((), xp, dtype) + return xp.flip(x) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flip_empty_dim_1(self, xp, dtype): + x = xp.array([], dtype).reshape((0,)) + return xp.flip(x, 0) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flip_empty_dim_2(self, xp, dtype): + x = xp.array([], dtype).reshape((0, 0)) + return xp.flip(x, 1) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flip_empty_dim_3(self, xp, dtype): + x = xp.array([], dtype).reshape((1, 0, 1)) + return xp.flip(x, 1) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_flip_empty_dim_all(self, xp, dtype): + x = xp.array([], dtype).reshape((1, 0, 1)) + return xp.flip(x) + + @testing.for_all_dtypes() + def test_flip_insufficient_ndim(self, dtype): + for xp in (numpy, cupy): + x = testing.shaped_arange((), xp, dtype) + with pytest.raises(ValueError): + xp.flip(x, 0) + + @testing.for_all_dtypes() + def test_flip_invalid_axis(self, dtype): + for xp in (numpy, cupy): + x = testing.shaped_arange((3, 4), xp, dtype) + with pytest.raises(ValueError): + xp.flip(x, 2) + + @testing.for_all_dtypes() + def test_flip_invalid_negative_axis(self, dtype): + for xp in (numpy, cupy): + x = testing.shaped_arange((3, 4), xp, dtype) + with pytest.raises(ValueError): + xp.flip(x, -3) + + +@pytest.mark.skip("`rot90` isn't supported yet") +class TestRot90(unittest.TestCase): + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_rot90_none(self, xp, dtype): + x = testing.shaped_arange((3, 4), xp, dtype) + return xp.rot90(x, 0) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_rot90_twice(self, xp, dtype): + x = testing.shaped_arange((3, 4, 2), xp, dtype) + return xp.rot90(x, 2) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_rot90_negative(self, xp, dtype): + x = testing.shaped_arange((3, 4, 2), xp, dtype) + return xp.rot90(x, -1) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_rot90_with_axes(self, xp, dtype): + x = testing.shaped_arange((3, 4, 2), xp, dtype) + return xp.rot90(x, 1, axes=(1, 2)) + + @testing.for_all_dtypes() + @testing.numpy_cupy_array_equal() + def test_rot90_with_negative_axes(self, xp, dtype): + x = testing.shaped_arange((3, 4, 2), xp, dtype) + return xp.rot90(x, 1, axes=(1, -1)) + + @testing.for_all_dtypes() + def test_rot90_insufficient_ndim(self, dtype): + for xp in (numpy, cupy): + x = testing.shaped_arange((3,), xp, dtype) + with pytest.raises(ValueError): + xp.rot90(x) + + @testing.for_all_dtypes() + def test_rot90_too_much_axes(self, dtype): + for xp in (numpy, cupy): + x = testing.shaped_arange((3, 4, 2), xp, dtype) + with pytest.raises(ValueError): + xp.rot90(x, 1, axes=(0, 1, 2)) + + @testing.for_all_dtypes() + def test_rot90_invalid_axes(self, dtype): + for xp in (numpy, cupy): + x = testing.shaped_arange((3, 4, 2), xp, dtype) + with pytest.raises(ValueError): + xp.rot90(x, 1, axes=(1, 3)) + + @testing.for_all_dtypes() + def test_rot90_invalid_negative_axes(self, dtype): + for xp in (numpy, cupy): + x = testing.shaped_arange((3, 4, 2), xp, dtype) + with pytest.raises(ValueError): + xp.rot90(x, 1, axes=(1, -2)) diff --git a/tests/third_party/cupy/manipulation_tests/test_transpose.py b/tests/third_party/cupy/manipulation_tests/test_transpose.py index ad71ad41b4b2..d758d15d3f77 100644 --- a/tests/third_party/cupy/manipulation_tests/test_transpose.py +++ b/tests/third_party/cupy/manipulation_tests/test_transpose.py @@ -103,7 +103,6 @@ def test_rollaxis(self, xp): a = testing.shaped_arange((2, 3, 4), xp) return xp.rollaxis(a, 2) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_rollaxis_failure(self): for xp in (numpy, cupy): a = testing.shaped_arange((2, 3, 4), xp) @@ -115,7 +114,6 @@ def test_swapaxes(self, xp): a = testing.shaped_arange((2, 3, 4), xp) return xp.swapaxes(a, 2, 0) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_swapaxes_failure(self): for xp in (numpy, cupy): a = testing.shaped_arange((2, 3, 4), xp) diff --git a/tests/third_party/cupy/math_tests/test_arithmetic.py b/tests/third_party/cupy/math_tests/test_arithmetic.py index f15682b9a675..c70c37e48a29 100644 --- a/tests/third_party/cupy/math_tests/test_arithmetic.py +++ b/tests/third_party/cupy/math_tests/test_arithmetic.py @@ -9,7 +9,7 @@ from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing -float_types = [numpy.float32, numpy.float64] +float_types = list(testing.helper._float_dtypes) complex_types = [] signed_int_types = [numpy.int32, numpy.int64] unsigned_int_types = [] @@ -126,7 +126,7 @@ def test_unary(self, xp): class ArithmeticBinaryBase: - @testing.numpy_cupy_allclose(atol=1e-4, type_check=has_support_aspect64()) + @testing.numpy_cupy_allclose(atol=1e-4, type_check=False) def check_binary(self, xp): arg1 = self.arg1 arg2 = self.arg2 @@ -135,20 +135,6 @@ def check_binary(self, xp): dtype1 = np1.dtype dtype2 = np2.dtype - if self.name == "power": - # TODO(niboshi): Fix this: power(0, 1j) - # numpy => 1+0j - # cupy => 0j - if dtype2 in complex_types and (np1 == 0).any(): - return xp.array(True) - - # TODO(niboshi): Fix this: xp.power(0j, 0) - # numpy => 1+0j - # cupy => 0j - c_arg1 = dtype1 in complex_types - if c_arg1 and (np1 == 0j).any() and (np2 == 0).any(): - return xp.array(True) - # TODO(niboshi): Fix this: xp.add(0j, xp.array([2.], 'f')).dtype # numpy => complex64 # cupy => complex128 @@ -186,36 +172,13 @@ def check_binary(self, xp): y = y.astype(numpy.complex64) # NumPy returns an output array of another type than DPNP when input ones have diffrent types. - if xp is cupy and dtype1 != dtype2 and not self.use_dtype: + if xp is numpy and dtype1 != dtype2: is_array_arg1 = not xp.isscalar(arg1) is_array_arg2 = not xp.isscalar(arg2) is_int_float = lambda _x, _y: numpy.issubdtype( _x, numpy.integer ) and numpy.issubdtype(_y, numpy.floating) - is_same_type = lambda _x, _y, _type: numpy.issubdtype( - _x, _type - ) and numpy.issubdtype(_y, _type) - - if self.name == "power": - if is_array_arg1 and is_array_arg2: - # If both inputs are arrays where one is of floating type and another - integer, - # NumPy will return an output array of always "float64" type, - # while DPNP will return the array of a wider type from the input arrays. - if is_int_float(dtype1, dtype2) or is_int_float( - dtype2, dtype1 - ): - y = y.astype(numpy.float64) - elif is_same_type( - dtype1, dtype2, numpy.floating - ) or is_same_type(dtype1, dtype2, numpy.integer): - # If one input is an array and another - scalar, - # NumPy will return an output array of the same type as the inpupt array has, - # while DPNP will return the array of a wider type from the inputs (considering both array and scalar). - if is_array_arg1 and not is_array_arg2: - y = y.astype(dtype1) - elif is_array_arg2 and not is_array_arg1: - y = y.astype(dtype2) return y diff --git a/tests/third_party/cupy/math_tests/test_hyperbolic.py b/tests/third_party/cupy/math_tests/test_hyperbolic.py index 6fa732c0200b..7d7fa6a6b13d 100644 --- a/tests/third_party/cupy/math_tests/test_hyperbolic.py +++ b/tests/third_party/cupy/math_tests/test_hyperbolic.py @@ -1,17 +1,17 @@ import unittest +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing -@testing.gpu class TestHyperbolic(unittest.TestCase): @testing.for_all_dtypes() - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) - @testing.for_dtypes(["f", "d"]) + @testing.for_dtypes(["e", "f", "d"]) @testing.numpy_cupy_allclose(atol=1e-5) def check_unary_unit(self, name, xp, dtype): a = xp.array([0.2, 0.4, 0.6, 0.8], dtype=dtype) @@ -29,7 +29,7 @@ def test_tanh(self): def test_arcsinh(self): self.check_unary("arcsinh") - @testing.for_dtypes(["f", "d"]) + @testing.for_dtypes(["e", "f", "d"]) @testing.numpy_cupy_allclose(atol=1e-5) def test_arccosh(self, xp, dtype): a = xp.array([1, 2, 3], dtype=dtype) diff --git a/tests/third_party/cupy/math_tests/test_rounding.py b/tests/third_party/cupy/math_tests/test_rounding.py index e17a48aa9e31..5c9cbcd4aa68 100644 --- a/tests/third_party/cupy/math_tests/test_rounding.py +++ b/tests/third_party/cupy/math_tests/test_rounding.py @@ -4,6 +4,7 @@ import pytest import dpnp as cupy +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing @@ -29,7 +30,7 @@ def check_unary_complex_unsupported(self, name, dtype): getattr(xp, name)(a) @testing.for_dtypes(["?", "b", "h", "i", "q", "e", "f", "d"]) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(type_check=False, atol=1e-5) def check_unary_negative(self, name, xp, dtype): a = xp.array([-3, -2, -1, 1, 2, 3], dtype=dtype) return getattr(xp, name)(a) @@ -70,8 +71,8 @@ def test_around(self): self.check_unary("around") self.check_unary_complex("around") - def test_round_(self): - self.check_unary("round_") + def test_round(self): + self.check_unary("round") self.check_unary_complex("around") @@ -102,7 +103,8 @@ def test_round(self, xp, dtype): @testing.numpy_cupy_array_equal() def test_round_out(self, xp): - a = testing.shaped_random(self.shape, xp, scale=100, dtype="d") + dtype = "d" if has_support_aspect64() else "f" + a = testing.shaped_random(self.shape, xp, scale=100, dtype=dtype) out = xp.empty_like(a) xp.around(a, self.decimals, out) return out @@ -115,16 +117,23 @@ def test_round_out(self, xp): } ) ) +@pytest.mark.usefixtures("allow_fall_back_on_numpy") class TestRoundExtreme(unittest.TestCase): shape = (20,) - @testing.for_dtypes([numpy.float64, numpy.complex128]) + dtype_ = ( + [numpy.float64, numpy.complex128] + if has_support_aspect64() + else [numpy.float32, numpy.complex64] + ) + + @testing.for_dtypes(dtype_) @testing.numpy_cupy_allclose() def test_round_large(self, xp, dtype): a = testing.shaped_random(self.shape, xp, scale=1e100, dtype=dtype) return xp.around(a, self.decimals) - @testing.for_dtypes([numpy.float64, numpy.complex128]) + @testing.for_dtypes(dtype_) @testing.numpy_cupy_allclose() def test_round_small(self, xp, dtype): a = testing.shaped_random(self.shape, xp, scale=1e-100, dtype=dtype) @@ -150,23 +159,23 @@ def test_round_small(self, xp, dtype): ) @pytest.mark.usefixtures("allow_fall_back_on_numpy") class TestRoundBorder(unittest.TestCase): - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_around_positive1(self, xp): a, decimals = self.value return xp.around(a, decimals) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_around_positive2(self, xp): a, decimals = self.value a = xp.asarray(a) return xp.around(a, decimals) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_around_negative1(self, xp): a, decimals = self.value return xp.around(-a, decimals) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_around_negative2(self, xp): a, decimals = self.value a = xp.asarray(a) diff --git a/tests/third_party/cupy/math_tests/test_trigonometric.py b/tests/third_party/cupy/math_tests/test_trigonometric.py index 60533205d930..778dc461454d 100644 --- a/tests/third_party/cupy/math_tests/test_trigonometric.py +++ b/tests/third_party/cupy/math_tests/test_trigonometric.py @@ -1,23 +1,24 @@ import unittest +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing class TestTrigonometric(unittest.TestCase): @testing.for_all_dtypes(no_complex=True) - @testing.numpy_cupy_allclose(atol=1e-5, type_check=False) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) @testing.for_all_dtypes(no_complex=True) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def check_binary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) b = testing.shaped_reverse_arange((2, 3), xp, dtype) return getattr(xp, name)(a, b) - @testing.for_dtypes(["f", "d"]) + @testing.for_dtypes(["e", "f", "d"]) @testing.numpy_cupy_allclose(atol=1e-5) def check_unary_unit(self, name, xp, dtype): a = xp.array([0.2, 0.4, 0.6, 0.8], dtype=dtype) diff --git a/tests/third_party/cupy/sorting_tests/test_search.py b/tests/third_party/cupy/sorting_tests/test_search.py index b7adc087632b..1a917367266d 100644 --- a/tests/third_party/cupy/sorting_tests/test_search.py +++ b/tests/third_party/cupy/sorting_tests/test_search.py @@ -694,8 +694,8 @@ def test_invalid_sorter(self): def test_nonint_sorter(self): for xp in (numpy, cupy): - x = testing.shaped_arange((12,), xp, xp.float64) + x = testing.shaped_arange((12,), xp, xp.float32) bins = xp.array([10, 4, 2, 1, 8]) - sorter = xp.array([], dtype=xp.float64) + sorter = xp.array([], dtype=xp.float32) with pytest.raises(TypeError): xp.searchsorted(bins, x, sorter=sorter) diff --git a/tests/third_party/cupy/statistics_tests/test_correlation.py b/tests/third_party/cupy/statistics_tests/test_correlation.py index 2353234077f2..406bf8b8a970 100644 --- a/tests/third_party/cupy/statistics_tests/test_correlation.py +++ b/tests/third_party/cupy/statistics_tests/test_correlation.py @@ -6,6 +6,7 @@ from dpctl import select_default_device import dpnp as cupy +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing @@ -181,7 +182,7 @@ def test_correlate_large_non_contiguous(self, xp, dtype): return xp.correlate(a[200::], b[10::700], mode=self.mode) @testing.for_all_dtypes_combination(names=["dtype1", "dtype2"]) - @testing.numpy_cupy_allclose(rtol=1e-2) + @testing.numpy_cupy_allclose(rtol=1e-2, type_check=has_support_aspect64()) def test_correlate_diff_types(self, xp, dtype1, dtype2): a = testing.shaped_random((200,), xp, dtype1) b = testing.shaped_random((100,), xp, dtype2) diff --git a/tests/third_party/cupy/statistics_tests/test_meanvar.py b/tests/third_party/cupy/statistics_tests/test_meanvar.py index 6b818237cfed..738057a99f49 100644 --- a/tests/third_party/cupy/statistics_tests/test_meanvar.py +++ b/tests/third_party/cupy/statistics_tests/test_meanvar.py @@ -4,6 +4,7 @@ import pytest import dpnp as cupy +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing ignore_runtime_warnings = pytest.mark.filterwarnings( @@ -14,21 +15,21 @@ @testing.gpu class TestMedian(unittest.TestCase): @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_median_noaxis(self, xp, dtype): a = testing.shaped_random((3, 4, 5), xp, dtype) return xp.median(a) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_median_axis1(self, xp, dtype): a = testing.shaped_random((3, 4, 5), xp, dtype) return xp.median(a, axis=1) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_median_axis2(self, xp, dtype): a = testing.shaped_random((3, 4, 5), xp, dtype) return xp.median(a, axis=2) @@ -42,14 +43,14 @@ def test_median_overwrite_input(self, xp, dtype): @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_median_keepdims_axis1(self, xp, dtype): a = testing.shaped_random((3, 4, 5), xp, dtype) return xp.median(a, axis=1, keepdims=True) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_median_keepdims_noaxis(self, xp, dtype): a = testing.shaped_random((3, 4, 5), xp, dtype) return xp.median(a, keepdims=True) @@ -91,7 +92,7 @@ def test_median_invalid_axis(self): @testing.gpu class TestMedianAxis(unittest.TestCase): @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_median_axis_sequence(self, xp, dtype): a = testing.shaped_random(self.shape, xp, dtype) return xp.median(a, self.axis, keepdims=self.keepdims) @@ -109,7 +110,7 @@ def test_average_all(self, xp, dtype): @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_average_axis(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return xp.average(a, axis=1) @@ -124,7 +125,7 @@ def test_average_weights(self, xp, dtype): @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_average_axis_weights(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) w = testing.shaped_arange((2, 3, 4), xp, dtype) @@ -154,25 +155,25 @@ def test_returned(self, dtype): @testing.gpu class TestMeanVar(unittest.TestCase): @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_mean_all(self, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return a.mean() @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_external_mean_all(self, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return xp.mean(a) @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_mean_axis(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return a.mean(axis=1) @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_external_mean_axis(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return xp.mean(a, axis=1) @@ -221,28 +222,28 @@ def test_external_var_all_ddof(self, xp, dtype): @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_var_axis(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return a.var(axis=1) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_external_var_axis(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return xp.var(a, axis=1) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_var_axis_ddof(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return a.var(axis=1, ddof=1) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_external_var_axis_ddof(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return xp.var(a, axis=1, ddof=1) @@ -273,28 +274,28 @@ def test_external_std_all_ddof(self, xp, dtype): @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_std_axis(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return a.std(axis=1) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_external_std_axis(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return xp.std(a, axis=1) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_std_axis_ddof(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return a.std(axis=1, ddof=1) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes() - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_external_std_axis_ddof(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) return xp.std(a, axis=1, ddof=1) @@ -490,7 +491,7 @@ def test_nanstd_float16(self, xp): @testing.gpu class TestProductZeroLength(unittest.TestCase): @testing.for_all_dtypes(no_complex=True) - @testing.numpy_cupy_allclose() + @testing.numpy_cupy_allclose(type_check=has_support_aspect64()) def test_external_mean_zero_len(self, xp, dtype): shape, axis = self.params a = testing.shaped_arange(shape, xp, dtype) diff --git a/tests/third_party/cupy/testing/helper.py b/tests/third_party/cupy/testing/helper.py index f9f419e91c37..9cf8403cdf2d 100644 --- a/tests/third_party/cupy/testing/helper.py +++ b/tests/third_party/cupy/testing/helper.py @@ -1005,6 +1005,18 @@ def decorator(impl): @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: + if ( + numpy.dtype(dtype).type in (numpy.float64, numpy.complex128) + and not has_support_aspect64() + ): + continue + + if ( + numpy.dtype(dtype).type == numpy.float16 + and not select_default_device().has_aspect_fp16 + ): + continue + try: kw[name] = numpy.dtype(dtype).type impl(*args, **kw) @@ -1019,15 +1031,19 @@ def test_func(*args, **kw): return decorator +def has_support_aspect64(): + return select_default_device().has_aspect_fp64 + + def _get_supported_float_dtypes(): - if select_default_device().has_aspect_fp64: + if has_support_aspect64(): return (numpy.float64, numpy.float32) else: return (numpy.float32,) def _get_supported_complex_dtypes(): - if select_default_device().has_aspect_fp64: + if has_support_aspect64(): return (numpy.complex128, numpy.complex64) else: return (numpy.complex64,)