implement dpnp.fft.fft and dpnp.fft.ifft using pybind11 extension

doc/known_words.txt

@@ -1,3 +1,4 @@
+al
 backend
 bitwise
 boolean
@@ -9,7 +10,9 @@ combinatorially
 conda
 cubically
 Decompositions
+diag
 dimensionality
+discretized
 docstring
 dpctl
 dpnp
@@ -19,9 +22,11 @@ einsum
 endian
 eps
 epsneg
+et
 Extrema
 finfo
 finiteness
+Flannery
 Fortran
 Frobenius
 Hadamard
@@ -41,10 +46,12 @@ ndarray
 ndarrays
 ndim
 normed
+Nyquist
 oneAPI
 orthonormal
 Penrose
 Polyutils
+pre
 prepend
 prepending
 representable
@@ -58,14 +65,21 @@ subclasses
 subtype
 SyclDevice
 SyclQueue
+tensordot
+Teukolsky
 th
+tril
+triu
+Tukey
 ufunc
 ufuncs
 Unary
+unscaled
 unicode
 usm
 Vandermonde
 vectorized
+Vetterline
 von
 Weibull
 whitespace

doc/reference/fft.rst

@@ -63,3 +63,5 @@ Helper routines
    .. dpnp.fft.config.set_cufft_gpus
    .. dpnp.fft.config.get_plan_cache
    .. dpnp.fft.config.show_plan_cache_info
+
+.. automodule:: dpnp.fft

dpnp/CMakeLists.txt

@@ -54,6 +54,7 @@ endfunction()
 
 add_subdirectory(backend)
 add_subdirectory(backend/extensions/blas)
+add_subdirectory(backend/extensions/fft)
 add_subdirectory(backend/extensions/lapack)
 add_subdirectory(backend/extensions/vm)
 add_subdirectory(backend/extensions/sycl_ext)

dpnp/backend/extensions/fft/CMakeLists.txt

@@ -0,0 +1,75 @@
+# *****************************************************************************
+# Copyright (c) 2024, 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.
+# *****************************************************************************
+
+
+set(python_module_name _fft_impl)
+set(_module_src
+    ${CMAKE_CURRENT_SOURCE_DIR}/fft_py.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/in_place.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/out_of_place.cpp
+)
+
+pybind11_add_module(${python_module_name} MODULE ${_module_src})
+add_sycl_to_target(TARGET ${python_module_name} SOURCES ${_module_src})
+
+if (WIN32)
+    if (${CMAKE_VERSION} VERSION_LESS "3.27")
+        # this is a work-around for target_link_options inserting option after -link option, cause
+        # linker to ignore it.
+        set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} -fsycl-device-code-split=per_kernel")
+    endif()
+endif()
+
+set_target_properties(${python_module_name} PROPERTIES CMAKE_POSITION_INDEPENDENT_CODE ON)
+
+target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../src)
+
+target_include_directories(${python_module_name} PUBLIC ${Dpctl_INCLUDE_DIRS})
+target_include_directories(${python_module_name} PUBLIC ${Dpctl_TENSOR_INCLUDE_DIR})
+
+if (WIN32)
+  target_compile_options(${python_module_name} PRIVATE
+    /clang:-fno-approx-func
+    /clang:-fno-finite-math-only
+    )
+else()
+  target_compile_options(${python_module_name} PRIVATE
+    -fno-approx-func
+    -fno-finite-math-only
+    )
+endif()
+
+target_link_options(${python_module_name} PUBLIC -fsycl-device-code-split=per_kernel)
+
+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_SYCL::DFT)
+
+install(TARGETS ${python_module_name}
+  DESTINATION "dpnp/backend/extensions/fft"
+)

dpnp/backend/extensions/fft/common.hpp

@@ -0,0 +1,239 @@
+//*****************************************************************************
+// Copyright (c) 2024, 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 <oneapi/mkl.hpp>
+#include <pybind11/pybind11.h>
+#include <sycl/sycl.hpp>
+
+namespace dpnp::extensions::fft
+{
+namespace mkl_dft = oneapi::mkl::dft;
+namespace py = pybind11;
+
+template <mkl_dft::precision prec, mkl_dft::domain dom>
+class DescriptorWrapper
+{
+public:
+    using descr_type = mkl_dft::descriptor<prec, dom>;
+
+    DescriptorWrapper(std::int64_t n) : descr_(n), queue_ptr_{} {}
+    DescriptorWrapper(std::vector<std::int64_t> dimensions)
+        : descr_(dimensions), queue_ptr_{}
+    {
+    }
+    ~DescriptorWrapper() {}
+
+    void commit(sycl::queue &q)
+    {
+        mkl_dft::precision fft_prec = get_precision();
+        if (fft_prec == mkl_dft::precision::DOUBLE &&
+            !q.get_device().has(sycl::aspect::fp64))
+        {
+            throw py::value_error("Descriptor is double precision but the "
+                                  "device does not support double precision.");
+        }
+
+        descr_.commit(q);
+        queue_ptr_ = std::make_unique<sycl::queue>(q);
+    }
+
+    descr_type &get_descriptor()
+    {
+        return descr_;
+    }
+
+    const sycl::queue &get_queue() const
+    {
+        if (queue_ptr_) {
+            return *queue_ptr_;
+        }
+        else {
+            throw std::runtime_error(
+                "Attempt to get queue when it is not yet set");
+        }
+    }
+
+    // config_param::DIMENSION
+    template <typename valT = std::int64_t>
+    const valT get_dim()
+    {
+        valT dim = -1;
+        descr_.get_value(mkl_dft::config_param::DIMENSION, &dim);
+
+        return dim;
+    }
+
+    // config_param::NUMBER_OF_TRANSFORMS
+    template <typename valT = std::int64_t>
+    const valT get_number_of_transforms()
+    {
+        valT transforms_count{};
+
+        descr_.get_value(mkl_dft::config_param::NUMBER_OF_TRANSFORMS,
+                         &transforms_count);
+        return transforms_count;
+    }
+
+    template <typename valT = std::int64_t>
+    void set_number_of_transforms(const valT &num)
+    {
+        descr_.set_value(mkl_dft::config_param::NUMBER_OF_TRANSFORMS, num);
+    }
+
+    // config_param::FWD_STRIDES
+    template <typename valT = std::vector<std::int64_t>>
+    const valT get_fwd_strides()
+    {
+        const typename valT::value_type dim = get_dim();
+
+        valT fwd_strides(dim + 1);
+        descr_.get_value(mkl_dft::config_param::FWD_STRIDES,
+                         fwd_strides.data());
+        return fwd_strides;
+    }
+
+    template <typename valT = std::vector<std::int64_t>>
+    void set_fwd_strides(const valT &strides)
+    {
+        const typename valT::value_type dim = get_dim();
+
+        if (static_cast<size_t>(dim + 1) != strides.size()) {
+            throw py::value_error(
+                "Strides length does not match descriptor's dimension");
+        }
+        descr_.set_value(mkl_dft::config_param::FWD_STRIDES, strides.data());
+    }
+
+    // config_param::BWD_STRIDES
+    template <typename valT = std::vector<std::int64_t>>
+    const valT get_bwd_strides()
+    {
+        const typename valT::value_type dim = get_dim();
+
+        valT bwd_strides(dim + 1);
+        descr_.get_value(mkl_dft::config_param::BWD_STRIDES,
+                         bwd_strides.data());
+        return bwd_strides;
+    }
+
+    template <typename valT = std::vector<std::int64_t>>
+    void set_bwd_strides(const valT &strides)
+    {
+        const typename valT::value_type dim = get_dim();
+
+        if (static_cast<size_t>(dim + 1) != strides.size()) {
+            throw py::value_error(
+                "Strides length does not match descriptor's dimension");
+        }
+        descr_.set_value(mkl_dft::config_param::BWD_STRIDES, strides.data());
+    }
+
+    // config_param::FWD_DISTANCE
+    template <typename valT = std::int64_t>
+    const valT get_fwd_distance()
+    {
+        valT dist = 0;
+
+        descr_.get_value(mkl_dft::config_param::FWD_DISTANCE, &dist);
+        return dist;
+    }
+
+    template <typename valT = std::int64_t>
+    void set_fwd_distance(const valT &dist)
+    {
+        descr_.set_value(mkl_dft::config_param::FWD_DISTANCE, dist);
+    }
+
+    // config_param::BWD_DISTANCE
+    template <typename valT = std::int64_t>
+    const valT get_bwd_distance()
+    {
+        valT dist = 0;
+
+        descr_.get_value(mkl_dft::config_param::BWD_DISTANCE, &dist);
+        return dist;
+    }
+
+    template <typename valT = std::int64_t>
+    void set_bwd_distance(const valT &dist)
+    {
+        descr_.set_value(mkl_dft::config_param::BWD_DISTANCE, dist);
+    }
+
+    // config_param::PLACEMENT
+    bool get_in_place()
+    {
+        // TODO: replace when MKLD-10506 is implemented
+        // mkl_dft::config_value placement;
+        DFTI_CONFIG_VALUE placement;
+
+        descr_.get_value(mkl_dft::config_param::PLACEMENT, &placement);
+        // TODO: replace when MKLD-10506 is implemented
+        // return (placement == mkl_dft::config_value::INPLACE);
+        return (placement == DFTI_CONFIG_VALUE::DFTI_INPLACE);
+    }
+
+    void set_in_place(const bool &in_place_request)
+    {
+        // TODO: replace when MKLD-10506 is implemented
+        // descr_.set_value(mkl_dft::config_param::PLACEMENT, (in_place_request)
+        // ? mkl_dft::config_value::INPLACE :
+        // mkl_dft::config_value::NOT_INPLACE);
+        descr_.set_value(mkl_dft::config_param::PLACEMENT,
+                         (in_place_request)
+                             ? DFTI_CONFIG_VALUE::DFTI_INPLACE
+                             : DFTI_CONFIG_VALUE::DFTI_NOT_INPLACE);
+    }
+
+    // config_param::PRECISION
+    mkl_dft::precision get_precision()
+    {
+        mkl_dft::precision fft_prec;
+
+        descr_.get_value(mkl_dft::config_param::PRECISION, &fft_prec);
+        return fft_prec;
+    }
+
+    // config_param::COMMIT_STATUS
+    bool is_committed()
+    {
+        // TODO: replace when MKLD-10506 is implemented
+        // mkl_dft::config_value committed;
+        DFTI_CONFIG_VALUE committed;
+
+        descr_.get_value(mkl_dft::config_param::COMMIT_STATUS, &committed);
+        // TODO: replace when MKLD-10506 is implemented
+        // return (committed == mkl_dft::config_value::COMMITTED);
+        return (committed == DFTI_CONFIG_VALUE::DFTI_COMMITTED);
+    }
+
+private:
+    mkl_dft::descriptor<prec, dom> descr_;
+    std::unique_ptr<sycl::queue> queue_ptr_;
+};
+
+} // namespace dpnp::extensions::fft