Add FFT LinOp

INSTALL.md

@@ -167,6 +167,8 @@ imposed by the `HIP` tool suite. The variables are the following:
   `hipBLAS` installation path. The default value is `${ROCM_PATH}/hipblas`.
 + CMake `-DHIPSPARSE_PATH=` or environment `export HIPSPARSE_PATH=`: sets the
   `hipSPARSE` installation path. The default value is `${ROCM_PATH}/hipsparse`.
++ CMake `-DHIPFFT_PATH=` or environment `export HIPFFT_PATH=`: sets the
+  `hipFFT` installation path. The default value is `${ROCM_PATH}/hipfft`.
 + CMake `-DROCRAND_PATH=` or environment `export ROCRAND_PATH=`: sets the
   `rocRAND` installation path. The default value is `${ROCM_PATH}/rocrand`.
 + CMake `-DHIPRAND_PATH=` or environment `export HIPRAND_PATH=`: sets the

README.md

@@ -57,6 +57,7 @@ The Ginkgo HIP module has the following __additional__ requirements:
 *    the HIP, hipBLAS, hipSPARSE, hip/rocRAND and rocThrust packages compiled with either:
     * _AMD_ backend (using the `clang` compiler)
     * _9.2 <= CUDA < 11_ backend
+* if the hipFFT package is available, it is used to implement the FFT LinOps.
 
 The Ginkgo DPC++ module has the following __additional__ requirements:
 

cmake/GinkgoConfig.cmake.in

@@ -165,6 +165,7 @@ endif()
 if((NOT GINKGO_BUILD_SHARED_LIBS) AND GINKGO_BUILD_HIP)
     find_package(HIP REQUIRED)
     find_package(hipblas REQUIRED)
+    find_package(hipfft) # optional
     find_package(hiprand REQUIRED)
     find_package(hipsparse REQUIRED)
     find_package(rocrand REQUIRED)

cmake/create_test.cmake

@@ -131,6 +131,7 @@ ginkgo_build_test_name(${test_name} test_target_name)
         # Only `exception_helpers` requires these so far, but it's much easier
         # to put these this way.
         ${HIPBLAS_INCLUDE_DIRS}
+        ${HIPFFT_INCLUDE_DIRS}
         ${hiprand_INCLUDE_DIRS}
         ${HIPSPARSE_INCLUDE_DIRS}
         )

core/CMakeLists.txt

@@ -26,6 +26,7 @@ target_sources(ginkgo
     matrix/diagonal.cpp
     matrix/ell.cpp
     matrix/fbcsr.cpp
+    matrix/fft.cpp
     matrix/hybrid.cpp
     matrix/identity.cpp
     matrix/permutation.cpp

core/base/allocator.hpp

@@ -73,25 +73,65 @@ class ExecutorAllocator {
      * @tparam ExecType  the static type of the executor
      */
     template <typename ExecType>
-    ExecutorAllocator(std::shared_ptr<ExecType> exec) : exec_{std::move(exec)}
+    ExecutorAllocator(std::shared_ptr<ExecType> exec) noexcept
+        : exec_{std::move(exec)}
     {}
 
+    /**
+     * Copy-constructs an allocator.
+     *
+     * @param other  the other allocator
+     */
+    ExecutorAllocator(const ExecutorAllocator& other) noexcept
+        : exec_{other.get_executor()}
+    {}
+
+    /**
+     * Copy-assigns an allocator.
+     *
+     * @param other  the other allocator
+     */
+    ExecutorAllocator& operator=(const ExecutorAllocator& other) noexcept
+    {
+        exec_ = other.get_executor();
+        return *this;
+    }
+
+    /**
+     * Copy-assigns an allocator.
+     *
+     * This is related to `std::allocator_traits::template rebind<U>` and its
+     * use in more advanced data structures.
+     *
+     * @param other  the other allocator
+     * @tparam U  the element type of the allocator to be assigned.
+     */
+    template <typename U>
+    ExecutorAllocator& operator=(const ExecutorAllocator<U>& other) noexcept
+    {
+        exec_ = other.get_executor();
+        return *this;
+    }
+
     /**
      * Constructs an allocator for another element type from a given executor.
      *
      * This is related to `std::allocator_traits::template rebind<U>` and its
      * use in more advanced data structures.
      *
-     * @param other  the other executor
+     * @param other  the other allocator
      * @tparam U  the element type of the allocator to be constructed.
      */
     template <typename U>
-    ExecutorAllocator(const ExecutorAllocator<U>& other)
+    ExecutorAllocator(const ExecutorAllocator<U>& other) noexcept
         : exec_{other.get_executor()}
     {}
 
     /** Returns the executor used by this allocator.  */
-    std::shared_ptr<const Executor> get_executor() const { return exec_; }
+    std::shared_ptr<const Executor> get_executor() const noexcept
+    {
+        return exec_;
+    }
 
     /**
      * Allocates a memory area of the given size.
@@ -119,7 +159,7 @@ class ExecutorAllocator {
      */
     template <typename T2>
     friend bool operator==(const ExecutorAllocator<T>& l,
-                           const ExecutorAllocator<T2>& r)
+                           const ExecutorAllocator<T2>& r) noexcept
     {
         return l.get_executor() == r.get_executor();
     }
@@ -133,7 +173,7 @@ class ExecutorAllocator {
      */
     template <typename T2>
     friend bool operator!=(const ExecutorAllocator<T>& l,
-                           const ExecutorAllocator<T2>& r)
+                           const ExecutorAllocator<T2>& r) noexcept
     {
         return !(l == r);
     }

core/device_hooks/common_kernels.inc.cpp

@@ -51,6 +51,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "core/matrix/diagonal_kernels.hpp"
 #include "core/matrix/ell_kernels.hpp"
 #include "core/matrix/fbcsr_kernels.hpp"
+#include "core/matrix/fft_kernels.hpp"
 #include "core/matrix/hybrid_kernels.hpp"
 #include "core/matrix/sellp_kernels.hpp"
 #include "core/matrix/sparsity_csr_kernels.hpp"
@@ -1031,6 +1032,28 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
 }  // namespace ell
 
 
+namespace fft {
+
+
+template <typename ValueType>
+GKO_DECLARE_FFT_KERNEL(ValueType)
+GKO_NOT_COMPILED(GKO_HOOK_MODULE);
+GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_TYPE(GKO_DECLARE_FFT_KERNEL);
+
+template <typename ValueType>
+GKO_DECLARE_FFT2_KERNEL(ValueType)
+GKO_NOT_COMPILED(GKO_HOOK_MODULE);
+GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_TYPE(GKO_DECLARE_FFT2_KERNEL);
+
+template <typename ValueType>
+GKO_DECLARE_FFT3_KERNEL(ValueType)
+GKO_NOT_COMPILED(GKO_HOOK_MODULE);
+GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_TYPE(GKO_DECLARE_FFT3_KERNEL);
+
+
+}  // namespace fft
+
+
 namespace hybrid {
 
 

core/device_hooks/cuda_hooks.cpp

@@ -139,6 +139,12 @@ std::string CusparseError::get_error(int64)
 }
 
 
+std::string CufftError::get_error(int64)
+{
+    return "ginkgo CUDA module is not compiled";
+}
+
+
 int CudaExecutor::get_num_devices() { return 0; }
 
 

core/device_hooks/hip_hooks.cpp

@@ -136,6 +136,12 @@ std::string HipsparseError::get_error(int64)
 }
 
 
+std::string HipfftError::get_error(int64)
+{
+    return "ginkgo HIP module is not compiled";
+}
+
+
 int HipExecutor::get_num_devices() { return 0; }