diff --git a/common/src/KokkosFFT_normalization.hpp b/common/src/KokkosFFT_normalization.hpp index f3d3ace9..b82ef2c2 100644 --- a/common/src/KokkosFFT_normalization.hpp +++ b/common/src/KokkosFFT_normalization.hpp @@ -12,13 +12,13 @@ namespace KokkosFFT { ORTHO }; - template - void _normalize(ViewType& inout, const T coef) { + template + void _normalize(const ExecutionSpace& exec_space, ViewType& inout, const T coef) { std::size_t size = inout.size(); auto* data = inout.data(); Kokkos::parallel_for( - Kokkos::RangePolicy>{0, size}, + Kokkos::RangePolicy>(exec_space, 0, size), KOKKOS_LAMBDA(const int& i) { data[i] *= coef; } ); } @@ -53,10 +53,10 @@ namespace KokkosFFT { return std::tuple ({coef, to_normalize}); } - template - void normalize(ViewType& inout, FFTDirectionType direction, Normalization normalization, std::size_t fft_size) { + template + void normalize(const ExecutionSpace& exec_space, ViewType& inout, FFTDirectionType direction, Normalization normalization, std::size_t fft_size) { auto [coef, to_normalize] = _coefficients(inout, direction, normalization, fft_size); - if(to_normalize) _normalize(inout, coef); + if(to_normalize) _normalize(exec_space, inout, coef); } }; diff --git a/common/src/KokkosFFT_transpose.hpp b/common/src/KokkosFFT_transpose.hpp index e735d6f9..6ebfdd2f 100644 --- a/common/src/KokkosFFT_transpose.hpp +++ b/common/src/KokkosFFT_transpose.hpp @@ -79,13 +79,13 @@ namespace KokkosFFT { return get_map_axes(view, axis_type<1>({axis})); } - template = nullptr> - void _transpose(InViewType& in, OutViewType& out, [[maybe_unused]] axis_type<1> _map) { + void _transpose(const ExecutionSpace& exec_space, InViewType& in, OutViewType& out, [[maybe_unused]] axis_type<1> _map) { } - template - void _transpose(InViewType& in, OutViewType& out, axis_type<2> _map) { + template + void _transpose(const ExecutionSpace& exec_space, InViewType& in, OutViewType& out, axis_type<2> _map) { constexpr std::size_t DIM = 2; static_assert(InViewType::rank() == DIM, "KokkosFFT::_transpose: Rank of View must be equal to Rank of transpose axes."); @@ -93,7 +93,7 @@ namespace KokkosFFT { constexpr std::size_t rank = InViewType::rank(); using array_layout_type = typename InViewType::array_layout; - using range_type = Kokkos::MDRangePolicy< Kokkos::Rank >; + using range_type = Kokkos::MDRangePolicy >; using tile_type = typename range_type::tile_type; using point_type = typename range_type::point_type; @@ -126,15 +126,15 @@ namespace KokkosFFT { ); } - template - void _transpose(InViewType& in, OutViewType& out, axis_type<3> _map) { + template + void _transpose(const ExecutionSpace& exec_space, InViewType& in, OutViewType& out, axis_type<3> _map) { constexpr std::size_t DIM = 3; static_assert(InViewType::rank() == DIM, "KokkosFFT::_transpose: Rank of View must be equal to Rank of transpose axes."); constexpr std::size_t rank = InViewType::rank(); using array_layout_type = typename InViewType::array_layout; - using range_type = Kokkos::MDRangePolicy< Kokkos::Rank >; + using range_type = Kokkos::MDRangePolicy >; using tile_type = typename range_type::tile_type; using point_type = typename range_type::point_type; @@ -197,8 +197,8 @@ namespace KokkosFFT { D (5, 7) and axis = -1 -> D' (5, 7) * */ - template - void transpose(InViewType& in, OutViewType& out, axis_type _map) { + template + void transpose(const ExecutionSpace& exec_space, InViewType& in, OutViewType& out, axis_type _map) { using in_value_type = typename InViewType::non_const_value_type; using out_value_type = typename OutViewType::non_const_value_type; using array_layout_type = typename InViewType::array_layout; @@ -215,7 +215,7 @@ namespace KokkosFFT { throw std::runtime_error("KokkosFFT::transpose: transpose not necessary"); } - _transpose(in, out, _map); + _transpose(exec_space, in, out, _map); } }; diff --git a/common/unit_test/Test_Normalization.cpp b/common/unit_test/Test_Normalization.cpp index a479642c..3e5db8f5 100644 --- a/common/unit_test/Test_Normalization.cpp +++ b/common/unit_test/Test_Normalization.cpp @@ -1,4 +1,3 @@ - #include #include #include "KokkosFFT_normalization.hpp" @@ -21,11 +20,11 @@ TEST(Normalization, Forward) { Kokkos::fence(); // Backward FFT with Forward Normalization -> Do nothing - KokkosFFT::normalize(x, KOKKOS_FFT_BACKWARD, KokkosFFT::Normalization::FORWARD, len); + KokkosFFT::normalize(execution_space(), x, KOKKOS_FFT_BACKWARD, KokkosFFT::Normalization::FORWARD, len); EXPECT_TRUE( allclose(x, ref_b, 1.e-5, 1.e-12) ); // Forward FFT with Forward Normalization -> 1/N normalization - KokkosFFT::normalize(x, KOKKOS_FFT_FORWARD, KokkosFFT::Normalization::FORWARD, len); + KokkosFFT::normalize(execution_space(), x, KOKKOS_FFT_FORWARD, KokkosFFT::Normalization::FORWARD, len); EXPECT_TRUE( allclose(x, ref_f, 1.e-5, 1.e-12) ); } @@ -45,11 +44,11 @@ TEST(Normalization, Backward) { Kokkos::fence(); // Forward FFT with Backward Normalization -> Do nothing - KokkosFFT::normalize(x, KOKKOS_FFT_FORWARD, KokkosFFT::Normalization::BACKWARD, len); + KokkosFFT::normalize(execution_space(), x, KOKKOS_FFT_FORWARD, KokkosFFT::Normalization::BACKWARD, len); EXPECT_TRUE( allclose(x, ref_f, 1.e-5, 1.e-12) ); // Backward FFT with Backward Normalization -> 1/N normalization - KokkosFFT::normalize(x, KOKKOS_FFT_BACKWARD, KokkosFFT::Normalization::BACKWARD, len); + KokkosFFT::normalize(execution_space(), x, KOKKOS_FFT_BACKWARD, KokkosFFT::Normalization::BACKWARD, len); EXPECT_TRUE( allclose(x, ref_b, 1.e-5, 1.e-12) ); } @@ -72,10 +71,10 @@ TEST(Normalization, Ortho) { Kokkos::fence(); // Forward FFT with Ortho Normalization -> 1 / sqrt(N) normalization - KokkosFFT::normalize(x_f, KOKKOS_FFT_FORWARD, KokkosFFT::Normalization::ORTHO, len); + KokkosFFT::normalize(execution_space(), x_f, KOKKOS_FFT_FORWARD, KokkosFFT::Normalization::ORTHO, len); EXPECT_TRUE( allclose(x_f, ref_f, 1.e-5, 1.e-12) ); // Backward FFT with Ortho Normalization -> 1 / sqrt(N) normalization - KokkosFFT::normalize(x_b, KOKKOS_FFT_BACKWARD, KokkosFFT::Normalization::ORTHO, len); + KokkosFFT::normalize(execution_space(), x_b, KOKKOS_FFT_BACKWARD, KokkosFFT::Normalization::ORTHO, len); EXPECT_TRUE( allclose(x_b, ref_b, 1.e-5, 1.e-12) ); } \ No newline at end of file diff --git a/common/unit_test/Test_Transpose.cpp b/common/unit_test/Test_Transpose.cpp index af0a6a1f..4736d89b 100644 --- a/common/unit_test/Test_Transpose.cpp +++ b/common/unit_test/Test_Transpose.cpp @@ -289,7 +289,7 @@ void test_transpose_1d() { Kokkos::fence(); EXPECT_THROW( - KokkosFFT::transpose(x, xt, axes_type<1>({0})), + KokkosFFT::transpose(execution_space(), x, xt, axes_type<1>({0})), std::runtime_error ); } @@ -325,11 +325,11 @@ TEST(Transpose, 2DLeftView) { Kokkos::fence(); EXPECT_THROW( - KokkosFFT::transpose(x, xt_axis0, axes_type<2>({0, 1})), // xt is identical to x + KokkosFFT::transpose(execution_space(), x, xt_axis0, axes_type<2>({0, 1})), // xt is identical to x std::runtime_error ); - KokkosFFT::transpose(x, xt_axis1, axes_type<2>({1, 0})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis1, axes_type<2>({1, 0})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis1, ref_axis1, 1.e-5, 1.e-12) ); } @@ -354,11 +354,11 @@ TEST(Transpose, 2DRightView) { Kokkos::deep_copy(ref_axis0, h_ref_axis0); Kokkos::fence(); - KokkosFFT::transpose(x, xt_axis0, axes_type<2>({1, 0})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis0, axes_type<2>({1, 0})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis0, ref_axis0, 1.e-5, 1.e-12) ); EXPECT_THROW( - KokkosFFT::transpose(x, xt_axis1, axes_type<2>({0, 1})), // xt is identical to x + KokkosFFT::transpose(execution_space(), x, xt_axis1, axes_type<2>({0, 1})), // xt is identical to x std::runtime_error ); } @@ -404,23 +404,23 @@ TEST(Transpose, 3DLeftView) { Kokkos::fence(); EXPECT_THROW( - KokkosFFT::transpose(x, xt_axis012, axes_type<3>({0, 1, 2})), // xt is identical to x + KokkosFFT::transpose(execution_space(), x, xt_axis012, axes_type<3>({0, 1, 2})), // xt is identical to x std::runtime_error ); - KokkosFFT::transpose(x, xt_axis021, axes_type<3>({0, 2, 1})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis021, axes_type<3>({0, 2, 1})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis021, ref_axis021, 1.e-5, 1.e-12) ); - KokkosFFT::transpose(x, xt_axis102, axes_type<3>({1, 0, 2})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis102, axes_type<3>({1, 0, 2})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis102, ref_axis102, 1.e-5, 1.e-12) ); - KokkosFFT::transpose(x, xt_axis120, axes_type<3>({1, 2, 0})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis120, axes_type<3>({1, 2, 0})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis120, ref_axis120, 1.e-5, 1.e-12) ); - KokkosFFT::transpose(x, xt_axis201, axes_type<3>({2, 0, 1})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis201, axes_type<3>({2, 0, 1})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis201, ref_axis201, 1.e-5, 1.e-12) ); - KokkosFFT::transpose(x, xt_axis210, axes_type<3>({2, 1, 0})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis210, axes_type<3>({2, 1, 0})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis210, ref_axis210, 1.e-5, 1.e-12) ); } @@ -465,22 +465,22 @@ TEST(Transpose, 3DRightView) { Kokkos::fence(); EXPECT_THROW( - KokkosFFT::transpose(x, xt_axis012, axes_type<3>({0, 1, 2})), // xt is identical to x + KokkosFFT::transpose(execution_space(), x, xt_axis012, axes_type<3>({0, 1, 2})), // xt is identical to x std::runtime_error ); - KokkosFFT::transpose(x, xt_axis021, axes_type<3>({0, 2, 1})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis021, axes_type<3>({0, 2, 1})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis021, ref_axis021, 1.e-5, 1.e-12) ); - KokkosFFT::transpose(x, xt_axis102, axes_type<3>({1, 0, 2})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis102, axes_type<3>({1, 0, 2})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis102, ref_axis102, 1.e-5, 1.e-12) ); - KokkosFFT::transpose(x, xt_axis120, axes_type<3>({1, 2, 0})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis120, axes_type<3>({1, 2, 0})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis120, ref_axis120, 1.e-5, 1.e-12) ); - KokkosFFT::transpose(x, xt_axis201, axes_type<3>({2, 0, 1})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis201, axes_type<3>({2, 0, 1})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis201, ref_axis201, 1.e-5, 1.e-12) ); - KokkosFFT::transpose(x, xt_axis210, axes_type<3>({2, 1, 0})); // xt is the transpose of x + KokkosFFT::transpose(execution_space(), x, xt_axis210, axes_type<3>({2, 1, 0})); // xt is the transpose of x EXPECT_TRUE( allclose(xt_axis210, ref_axis210, 1.e-5, 1.e-12) ); } \ No newline at end of file diff --git a/examples/01_1DFFT/01_1DFFT.cpp b/examples/01_1DFFT/01_1DFFT.cpp index 3169f9eb..f9a7872a 100644 --- a/examples/01_1DFFT/01_1DFFT.cpp +++ b/examples/01_1DFFT/01_1DFFT.cpp @@ -20,22 +20,22 @@ int main( int argc, char* argv[] ) { Kokkos::Random_XorShift64_Pool<> random_pool(12345); Kokkos::fill_random(xc2c, random_pool, I); - KokkosFFT::fft(xc2c, xc2c_hat); - KokkosFFT::ifft(xc2c_hat, xc2c_inv); + KokkosFFT::fft(execution_space(), xc2c, xc2c_hat); + KokkosFFT::ifft(execution_space(), xc2c_hat, xc2c_inv); // 1D R2C FFT View1D xr2c("xr2c", n0); View1D > xr2c_hat("xr2c_hat", n0/2+1); Kokkos::fill_random(xr2c, random_pool, 1); - KokkosFFT::rfft(xr2c, xr2c_hat); + KokkosFFT::rfft(execution_space(), xr2c, xr2c_hat); // 1D C2R FFT View1D > xc2r("xr2c_hat", n0/2+1); View1D xc2r_hat("xc2r", n0); Kokkos::fill_random(xc2r, random_pool, I); - KokkosFFT::irfft(xc2r, xc2r_hat); + KokkosFFT::irfft(execution_space(), xc2r, xc2r_hat); } Kokkos::finalize(); diff --git a/examples/02_2DFFT/02_2DFFT.cpp b/examples/02_2DFFT/02_2DFFT.cpp index 0faeced7..5f837c85 100644 --- a/examples/02_2DFFT/02_2DFFT.cpp +++ b/examples/02_2DFFT/02_2DFFT.cpp @@ -20,22 +20,22 @@ int main( int argc, char* argv[] ) { Kokkos::Random_XorShift64_Pool<> random_pool(12345); Kokkos::fill_random(xc2c, random_pool, I); - KokkosFFT::fft2(xc2c, xc2c_hat); - KokkosFFT::ifft2(xc2c_hat, xc2c_inv); + KokkosFFT::fft2(execution_space(), xc2c, xc2c_hat); + KokkosFFT::ifft2(execution_space(), xc2c_hat, xc2c_inv); // 2D R2C FFT View2D xr2c("xr2c", n0, n1); View2D > xr2c_hat("xr2c_hat", n0, n1/2+1); Kokkos::fill_random(xr2c, random_pool, 1); - KokkosFFT::rfft2(xr2c, xr2c_hat); + KokkosFFT::rfft2(execution_space(), xr2c, xr2c_hat); // 2D C2R FFT View2D > xc2r("xr2c_hat", n0, n1/2+1); View2D xc2r_hat("xc2r", n0, n1); Kokkos::fill_random(xc2r, random_pool, I); - KokkosFFT::irfft2(xc2r, xc2r_hat); + KokkosFFT::irfft2(execution_space(), xc2r, xc2r_hat); } Kokkos::finalize(); diff --git a/examples/03_NDFFT/03_NDFFT.cpp b/examples/03_NDFFT/03_NDFFT.cpp index 991e67c3..4bb7b8be 100644 --- a/examples/03_NDFFT/03_NDFFT.cpp +++ b/examples/03_NDFFT/03_NDFFT.cpp @@ -20,22 +20,22 @@ int main( int argc, char* argv[] ) { Kokkos::Random_XorShift64_Pool<> random_pool(12345); Kokkos::fill_random(xc2c, random_pool, I); - KokkosFFT::fftn(xc2c, xc2c_hat); - KokkosFFT::ifftn(xc2c_hat, xc2c_inv); + KokkosFFT::fftn(execution_space(), xc2c, xc2c_hat); + KokkosFFT::ifftn(execution_space(), xc2c_hat, xc2c_inv); // 3D R2C FFT View3D xr2c("xr2c", n0, n1, n2); View3D > xr2c_hat("xr2c_hat", n0, n1, n2/2+1); Kokkos::fill_random(xr2c, random_pool, 1); - KokkosFFT::rfftn(xr2c, xr2c_hat); + KokkosFFT::rfftn(execution_space(), xr2c, xr2c_hat); // 3D C2R FFT View3D > xc2r("xr2c_hat", n0, n1, n2/2+1); View3D xc2r_hat("xc2r", n0, n1, n2); Kokkos::fill_random(xc2r, random_pool, I); - KokkosFFT::irfftn(xc2r, xc2r_hat); + KokkosFFT::irfftn(execution_space(), xc2r, xc2r_hat); } Kokkos::finalize(); diff --git a/examples/04_batchedFFT/04_batchedFFT.cpp b/examples/04_batchedFFT/04_batchedFFT.cpp index 9263d377..36e32a97 100644 --- a/examples/04_batchedFFT/04_batchedFFT.cpp +++ b/examples/04_batchedFFT/04_batchedFFT.cpp @@ -20,22 +20,22 @@ int main( int argc, char* argv[] ) { Kokkos::Random_XorShift64_Pool<> random_pool(12345); Kokkos::fill_random(xc2c, random_pool, I); - KokkosFFT::fft(xc2c, xc2c_hat, KokkosFFT::Normalization::BACKWARD, /*axis=*/-1); - KokkosFFT::ifft(xc2c_hat, xc2c_inv, KokkosFFT::Normalization::BACKWARD, /*axis=*/-1); + KokkosFFT::fft(execution_space(), xc2c, xc2c_hat, KokkosFFT::Normalization::BACKWARD, /*axis=*/-1); + KokkosFFT::ifft(execution_space(), xc2c_hat, xc2c_inv, KokkosFFT::Normalization::BACKWARD, /*axis=*/-1); // 1D batched R2C FFT View3D xr2c("xr2c", n0, n1, n2); View3D > xr2c_hat("xr2c_hat", n0, n1, n2/2+1); Kokkos::fill_random(xr2c, random_pool, 1); - KokkosFFT::rfft(xr2c, xr2c_hat, KokkosFFT::Normalization::BACKWARD, /*axis=*/-1); + KokkosFFT::rfft(execution_space(), xr2c, xr2c_hat, KokkosFFT::Normalization::BACKWARD, /*axis=*/-1); // 1D batched C2R FFT View3D > xc2r("xr2c_hat", n0, n1, n2/2+1); View3D xc2r_hat("xc2r", n0, n1, n2); Kokkos::fill_random(xc2r, random_pool, I); - KokkosFFT::irfft(xc2r, xc2r_hat, KokkosFFT::Normalization::BACKWARD, /*axis=*/-1); + KokkosFFT::irfft(execution_space(), xc2r, xc2r_hat, KokkosFFT::Normalization::BACKWARD, /*axis=*/-1); } Kokkos::finalize(); diff --git a/fft/src/KokkosFFT_Cuda_plans.hpp b/fft/src/KokkosFFT_Cuda_plans.hpp index dfc23896..8cf8a031 100644 --- a/fft/src/KokkosFFT_Cuda_plans.hpp +++ b/fft/src/KokkosFFT_Cuda_plans.hpp @@ -7,9 +7,9 @@ namespace KokkosFFT { // 1D transform - template = nullptr> - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -36,9 +36,9 @@ namespace KokkosFFT { } // 2D transform - template = nullptr> - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -63,9 +63,9 @@ namespace KokkosFFT { } // 3D transform - template = nullptr> - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -92,9 +92,9 @@ namespace KokkosFFT { } // ND transform - template = nullptr> - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -133,8 +133,8 @@ namespace KokkosFFT { } // batched transform, over ND Views - template - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction, axis_type axes) { + template + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction, axis_type axes) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, diff --git a/fft/src/KokkosFFT_HIP_plans.hpp b/fft/src/KokkosFFT_HIP_plans.hpp index bd2822b2..d0592ff3 100644 --- a/fft/src/KokkosFFT_HIP_plans.hpp +++ b/fft/src/KokkosFFT_HIP_plans.hpp @@ -6,10 +6,10 @@ #include "KokkosFFT_layouts.hpp" namespace KokkosFFT { - // 1D transform - template = nullptr> - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -28,24 +28,24 @@ namespace KokkosFFT { auto [in_extents, out_extents, fft_extents] = get_extents(in, out, axis); const int nx = fft_extents.at(0); int fft_size = std::accumulate(fft_extents.begin(), fft_extents.end(), 1, std::multiplies<>()); - + hipfft_rt = hipfftPlan1d(&plan, nx, type, batch); if(hipfft_rt != HIPFFT_SUCCESS) throw std::runtime_error("hipfftPlan1d failed"); return fft_size; } - // 2D transform - template = nullptr> - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, "KokkosFFT::_create: OutViewType is not a Kokkos::View."); using in_value_type = typename InViewType::non_const_value_type; using out_value_type = typename OutViewType::non_const_value_type; - + hipfftResult hipfft_rt = hipfftCreate(&plan); if(hipfft_rt != HIPFFT_SUCCESS) throw std::runtime_error("hipfftCreate failed"); @@ -62,10 +62,10 @@ namespace KokkosFFT { return fft_size; } - // 3D transform - template = nullptr> - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -82,7 +82,7 @@ namespace KokkosFFT { auto type = transform_type::type(); auto [in_extents, out_extents, fft_extents] = get_extents(in, out, axis); - + const int nx = fft_extents.at(0), ny = fft_extents.at(1), nz = fft_extents.at(2); int fft_size = std::accumulate(fft_extents.begin(), fft_extents.end(), 1, std::multiplies<>()); @@ -92,10 +92,10 @@ namespace KokkosFFT { return fft_size; } - // ND transform - template = nullptr> - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -115,9 +115,9 @@ namespace KokkosFFT { int idist = std::accumulate(in_extents.begin(), in_extents.end(), 1, std::multiplies<>()); int odist = std::accumulate(out_extents.begin(), out_extents.end(), 1, std::multiplies<>()); int fft_size = std::accumulate(fft_extents.begin(), fft_extents.end(), 1, std::multiplies<>()); - + hipfft_rt = hipfftPlanMany( - &plan, + &plan, rank, fft_extents.data(), nullptr, @@ -134,8 +134,8 @@ namespace KokkosFFT { } // batched transform, over ND Views - template - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction, axis_type axes) { + template + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType direction, axis_type axes) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -163,7 +163,7 @@ namespace KokkosFFT { throw std::runtime_error("hipfftCreate failed"); hipfft_rt = hipfftPlanMany( - &plan, + &plan, rank, fft_extents.data(), in_extents.data(), diff --git a/fft/src/KokkosFFT_OpenMP_plans.hpp b/fft/src/KokkosFFT_OpenMP_plans.hpp index 35323cc3..6ad5f0e4 100644 --- a/fft/src/KokkosFFT_OpenMP_plans.hpp +++ b/fft/src/KokkosFFT_OpenMP_plans.hpp @@ -6,26 +6,22 @@ #include "KokkosFFT_layouts.hpp" namespace KokkosFFT { - template - void _init_threads() { - #if defined(KOKKOS_ENABLE_OPENMP) - int nthreads=0; - #pragma omp parallel - nthreads = omp_get_num_threads(); + template + void _init_threads(const ExecutionSpace& exec_space) { + int nthreads = exec_space.concurrency(); - if constexpr (std::is_same_v) { - fftwf_init_threads(); - fftwf_plan_with_nthreads(nthreads); - } else { - fftw_init_threads(); - fftw_plan_with_nthreads(nthreads); - } - #endif + if constexpr (std::is_same_v) { + fftwf_init_threads(); + fftwf_plan_with_nthreads(nthreads); + } else { + fftw_init_threads(); + fftw_plan_with_nthreads(nthreads); + } } // ND transform - template - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType sign) { + template + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType sign) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -33,7 +29,7 @@ namespace KokkosFFT { using in_value_type = typename InViewType::non_const_value_type; using out_value_type = typename OutViewType::non_const_value_type; - _init_threads>(); + _init_threads>(exec_space); const int rank = InViewType::rank(); const int axis = -1; @@ -140,8 +136,8 @@ namespace KokkosFFT { } // batched transform, over ND Views - template - auto _create(PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType sign, axis_type axes) { + template + auto _create(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, const OutViewType& out, [[maybe_unused]] FFTDirectionType sign, axis_type axes) { static_assert(Kokkos::is_view::value, "KokkosFFT::_create: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -153,7 +149,7 @@ namespace KokkosFFT { "KokkosFFT::_create: Rank of View must be larger than Rank of FFT."); const int rank = fft_rank; - _init_threads>(); + _init_threads>(exec_space); constexpr auto type = transform_type::type(); auto [in_extents, out_extents, fft_extents, howmany] = get_extents_batched(in, out, axes); diff --git a/fft/src/KokkosFFT_Plans.hpp b/fft/src/KokkosFFT_Plans.hpp index 1f4d56fb..44ec4ad4 100644 --- a/fft/src/KokkosFFT_Plans.hpp +++ b/fft/src/KokkosFFT_Plans.hpp @@ -24,7 +24,7 @@ #endif namespace KokkosFFT { - template + template class Plan { using in_value_type = typename InViewType::non_const_value_type; using out_value_type = typename OutViewType::non_const_value_type; @@ -45,7 +45,7 @@ namespace KokkosFFT { nonConstOutViewType m_out_T; public: - explicit Plan(InViewType& in, OutViewType& out) : m_fft_size(1), m_is_transpose_needed(false) { + explicit Plan(const ExecutionSpace& exec_space, InViewType& in, OutViewType& out) : m_fft_size(1), m_is_transpose_needed(false) { // Available only for R2C or C2R // For C2C, direction should be given by an user static_assert(Kokkos::is_view::value, @@ -65,10 +65,10 @@ namespace KokkosFFT { throw std::runtime_error("direction not specified for Complex to Complex transform"); } - m_fft_size = _create(m_plan, in, out, direction); + m_fft_size = _create(exec_space, m_plan, in, out, direction); } - explicit Plan(InViewType& in, OutViewType& out, axis_type axes) : m_fft_size(1), m_is_transpose_needed(false) { + explicit Plan(const ExecutionSpace& exec_space, InViewType& in, OutViewType& out, axis_type axes) : m_fft_size(1), m_is_transpose_needed(false) { // Available only for R2C or C2R // For C2C, direction should be given by an user static_assert(Kokkos::is_view::value, @@ -88,20 +88,20 @@ namespace KokkosFFT { throw std::runtime_error("direction not specified for Complex to Complex transform"); } - m_fft_size = _create(m_plan, in, out, direction, axes); + m_fft_size = _create(exec_space, m_plan, in, out, direction, axes); } - explicit Plan(InViewType& in, OutViewType& out, FFTDirectionType direction) : m_fft_size(1), m_is_transpose_needed(false) { + explicit Plan(const ExecutionSpace& exec_space, InViewType& in, OutViewType& out, FFTDirectionType direction) : m_fft_size(1), m_is_transpose_needed(false) { static_assert(Kokkos::is_view::value, "KokkosFFT::Plan: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, "KokkosFFT::Plan: OutViewType is not a Kokkos::View."); /* Apply FFT over entire axes or along inner most directions */ - m_fft_size = _create(m_plan, in, out, direction); + m_fft_size = _create(exec_space, m_plan, in, out, direction); } - explicit Plan(InViewType& in, OutViewType& out, FFTDirectionType direction, int axis) : m_fft_size(1), m_is_transpose_needed(false) { + explicit Plan(const ExecutionSpace& exec_space, InViewType& in, OutViewType& out, FFTDirectionType direction, int axis) : m_fft_size(1), m_is_transpose_needed(false) { static_assert(Kokkos::is_view::value, "KokkosFFT::Plan: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -109,10 +109,10 @@ namespace KokkosFFT { std::tie(m_map, m_map_inv) = KokkosFFT::get_map_axes(in, axis); m_is_transpose_needed = KokkosFFT::is_transpose_needed(m_map); - m_fft_size = _create(m_plan, in, out, direction, axis_type<1>{axis}); + m_fft_size = _create(exec_space, m_plan, in, out, direction, axis_type<1>{axis}); } - explicit Plan(InViewType& in, OutViewType& out, FFTDirectionType direction, axis_type axes) : m_fft_size(1), m_is_transpose_needed(false) { + explicit Plan(const ExecutionSpace& exec_space, InViewType& in, OutViewType& out, FFTDirectionType direction, axis_type axes) : m_fft_size(1), m_is_transpose_needed(false) { static_assert(Kokkos::is_view::value, "KokkosFFT::Plan: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -120,7 +120,7 @@ namespace KokkosFFT { std::tie(m_map, m_map_inv) = KokkosFFT::get_map_axes(in, axes); m_is_transpose_needed = KokkosFFT::is_transpose_needed(m_map); - m_fft_size = _create(m_plan, in, out, direction, axes); + m_fft_size = _create(exec_space, m_plan, in, out, direction, axes); } ~Plan() { diff --git a/fft/src/KokkosFFT_Transform.hpp b/fft/src/KokkosFFT_Transform.hpp index 7a059801..6e78bc89 100644 --- a/fft/src/KokkosFFT_Transform.hpp +++ b/fft/src/KokkosFFT_Transform.hpp @@ -27,8 +27,8 @@ // 1D Transform namespace KokkosFFT { - template - void _fft(PlanType& plan, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { + template + void _fft(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { static_assert(Kokkos::is_view::value, "KokkosFFT::_fft: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -41,11 +41,11 @@ namespace KokkosFFT { auto* odata = reinterpret_cast::type*>(out.data()); _exec(plan.plan(), idata, odata, KOKKOS_FFT_FORWARD); - normalize(out, KOKKOS_FFT_FORWARD, norm, plan.fft_size()); + normalize(exec_space, out, KOKKOS_FFT_FORWARD, norm, plan.fft_size()); } - template - void _ifft(PlanType& plan, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { + template + void _ifft(const ExecutionSpace& exec_space, PlanType& plan, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { static_assert(Kokkos::is_view::value, "KokkosFFT::_ifft: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -58,59 +58,71 @@ namespace KokkosFFT { auto* odata = reinterpret_cast::type*>(out.data()); _exec(plan.plan(), idata, odata, KOKKOS_FFT_BACKWARD); - normalize(out, KOKKOS_FFT_BACKWARD, norm, plan.fft_size()); + normalize(exec_space, out, KOKKOS_FFT_BACKWARD, norm, plan.fft_size()); } - template - void fft(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, int axis=-1) { + template + void fft(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, int axis=-1) { static_assert(Kokkos::is_view::value, "KokkosFFT::fft: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, "KokkosFFT::fft: OutViewType is not a Kokkos::View."); - Plan plan(in, out, KOKKOS_FFT_FORWARD, axis); + static_assert( + Kokkos::SpaceAccessibility::accessible, + "KokkosFFT::fft: execution_space cannot access data in InViewType" + ); + + static_assert( + Kokkos::SpaceAccessibility::accessible, + "KokkosFFT::fft: execution_space cannot access data in OutViewType" + ); + + Plan plan(exec_space, in, out, KOKKOS_FFT_FORWARD, axis); if(plan.is_transpose_needed()) { InViewType in_T; OutViewType out_T; - KokkosFFT::transpose(in, in_T, plan.map()); - KokkosFFT::transpose(out, out_T, plan.map()); + KokkosFFT::transpose(exec_space, in, in_T, plan.map()); + KokkosFFT::transpose(exec_space, out, out_T, plan.map()); - _fft(plan, in_T, out_T, norm); + _fft(exec_space, plan, in_T, out_T, norm); - KokkosFFT::transpose(out_T, out, plan.map_inv()); + KokkosFFT::transpose(exec_space, out_T, out, plan.map_inv()); } else { - _fft(plan, in, out, norm); + _fft(exec_space, plan, in, out, norm); } } - template - void ifft(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, int axis=-1) { + template + void ifft(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, int axis=-1) { static_assert(Kokkos::is_view::value, "KokkosFFT::ifft: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, "KokkosFFT::ifft: OutViewType is not a Kokkos::View."); - Plan plan(in, out, KOKKOS_FFT_BACKWARD, axis); + Plan plan(exec_space, in, out, KOKKOS_FFT_BACKWARD, axis); if(plan.is_transpose_needed()) { InViewType in_T; OutViewType out_T; - KokkosFFT::transpose(in, in_T, plan.map()); - KokkosFFT::transpose(out, out_T, plan.map()); + KokkosFFT::transpose(exec_space, in, in_T, plan.map()); + KokkosFFT::transpose(exec_space, out, out_T, plan.map()); - _ifft(plan, in_T, out_T, norm); + _ifft(exec_space, plan, in_T, out_T, norm); - KokkosFFT::transpose(out_T, out, plan.map_inv()); + KokkosFFT::transpose(exec_space, out_T, out, plan.map_inv()); } else { - _ifft(plan, in, out, norm); + _ifft(exec_space, plan, in, out, norm); } } - template - void rfft(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, int axis=-1) { + template + void rfft(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, int axis=-1) { static_assert(Kokkos::is_view::value, "KokkosFFT::rfft: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -124,11 +136,11 @@ namespace KokkosFFT { static_assert(is_complex::value, "KokkosFFT::rfft: OutViewType must be complex"); - fft(in, out, norm, axis); + fft(exec_space, in, out, norm, axis); } - template - void irfft(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, int axis=-1) { + template + void irfft(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, int axis=-1) { static_assert(Kokkos::is_view::value, "KokkosFFT::irfft: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -142,59 +154,57 @@ namespace KokkosFFT { static_assert(std::is_floating_point::value, "KokkosFFT::irfft: OutViewType must be real"); - ifft(in, out, norm, axis); + ifft(exec_space, in, out, norm, axis); } -}; -namespace KokkosFFT { - template - void fft2(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, axis_type<2> axes={-2, -1}) { + template + void fft2(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, axis_type<2> axes={-2, -1}) { static_assert(Kokkos::is_view::value, "KokkosFFT::fft2: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, "KokkosFFT::fft2: OutViewType is not a Kokkos::View."); - Plan plan(in, out, KOKKOS_FFT_FORWARD, axes); + Plan plan(exec_space, in, out, KOKKOS_FFT_FORWARD, axes); if(plan.is_transpose_needed()) { InViewType in_T; OutViewType out_T; - KokkosFFT::transpose(in, in_T, plan.map()); - KokkosFFT::transpose(out, out_T, plan.map()); + KokkosFFT::transpose(exec_space, in, in_T, plan.map()); + KokkosFFT::transpose(exec_space, out, out_T, plan.map()); - _fft(plan, in_T, out_T, norm); + _fft(exec_space, plan, in_T, out_T, norm); - KokkosFFT::transpose(out_T, out, plan.map_inv()); + KokkosFFT::transpose(exec_space, out_T, out, plan.map_inv()); } else { - _fft(plan, in, out, norm); + _fft(exec_space, plan, in, out, norm); } } - template - void ifft2(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, axis_type<2> axes={-2, -1}) { + template + void ifft2(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, axis_type<2> axes={-2, -1}) { static_assert(Kokkos::is_view::value, "KokkosFFT::ifft2: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, "KokkosFFT::ifft2: OutViewType is not a Kokkos::View."); - Plan plan(in, out, KOKKOS_FFT_BACKWARD, axes); + Plan plan(exec_space, in, out, KOKKOS_FFT_BACKWARD, axes); if(plan.is_transpose_needed()) { InViewType in_T; OutViewType out_T; - KokkosFFT::transpose(in, in_T, plan.map()); - KokkosFFT::transpose(out, out_T, plan.map()); + KokkosFFT::transpose(exec_space, in, in_T, plan.map()); + KokkosFFT::transpose(exec_space, out, out_T, plan.map()); - _ifft(plan, in_T, out_T, norm); + _ifft(exec_space, plan, in_T, out_T, norm); - KokkosFFT::transpose(out_T, out, plan.map_inv()); + KokkosFFT::transpose(exec_space, out_T, out, plan.map_inv()); } else { - _ifft(plan, in, out, norm); + _ifft(exec_space, plan, in, out, norm); } } - template - void rfft2(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, axis_type<2> axes={-2, -1}) { + template + void rfft2(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, axis_type<2> axes={-2, -1}) { static_assert(Kokkos::is_view::value, "KokkosFFT::rfft2: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -208,11 +218,11 @@ namespace KokkosFFT { static_assert(is_complex::value, "KokkosFFT::rfft2: OutViewType must be complex"); - fft2(in, out, norm, axes); + fft2(exec_space, in, out, norm, axes); } - template - void irfft2(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, axis_type<2> axes={-2, -1}) { + template + void irfft2(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD, axis_type<2> axes={-2, -1}) { static_assert(Kokkos::is_view::value, "KokkosFFT::irfft2: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -226,13 +236,11 @@ namespace KokkosFFT { static_assert(std::is_floating_point::value, "KokkosFFT::irfft2: OutViewType must be real"); - ifft2(in, out, norm, axes); + ifft2(exec_space, in, out, norm, axes); } -} -namespace KokkosFFT { - template - void fftn(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { + template + void fftn(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { static_assert(Kokkos::is_view::value, "KokkosFFT::fftn: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -243,47 +251,47 @@ namespace KokkosFFT { constexpr int start = -static_cast(rank); axis_type axes = index_sequence(start); - Plan plan(in, out, KOKKOS_FFT_FORWARD, axes); + Plan plan(exec_space, in, out, KOKKOS_FFT_FORWARD, axes); if(plan.is_transpose_needed()) { InViewType in_T; OutViewType out_T; - KokkosFFT::transpose(in, in_T, plan.map()); - KokkosFFT::transpose(out, out_T, plan.map()); + KokkosFFT::transpose(exec_space, in, in_T, plan.map()); + KokkosFFT::transpose(exec_space, out, out_T, plan.map()); - _fft(plan, in_T, out_T, norm); + _fft(exec_space, plan, in_T, out_T, norm); - KokkosFFT::transpose(out_T, out, plan.map_inv()); + KokkosFFT::transpose(exec_space, out_T, out, plan.map_inv()); } else { - _fft(plan, in, out, norm); + _fft(exec_space, plan, in, out, norm); } } - template - void fftn(const InViewType& in, OutViewType& out, axis_type axes, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { + template + void fftn(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, axis_type axes, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { static_assert(Kokkos::is_view::value, "KokkosFFT::fftn: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, "KokkosFFT::fftn: OutViewType is not a Kokkos::View."); - Plan plan(in, out, KOKKOS_FFT_FORWARD, axes); + Plan plan(exec_space, in, out, KOKKOS_FFT_FORWARD, axes); if(plan.is_transpose_needed()) { InViewType in_T; OutViewType out_T; - KokkosFFT::transpose(in, in_T, plan.map()); - KokkosFFT::transpose(out, out_T, plan.map()); + KokkosFFT::transpose(exec_space, in, in_T, plan.map()); + KokkosFFT::transpose(exec_space, out, out_T, plan.map()); - _fft(plan, in_T, out_T, norm); + _fft(exec_space, plan, in_T, out_T, norm); - KokkosFFT::transpose(out_T, out, plan.map_inv()); + KokkosFFT::transpose(exec_space, out_T, out, plan.map_inv()); } else { - _fft(plan, in, out, norm); + _fft(exec_space, plan, in, out, norm); } } - template - void ifftn(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { + template + void ifftn(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { static_assert(Kokkos::is_view::value, "KokkosFFT::ifftn: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -294,47 +302,47 @@ namespace KokkosFFT { constexpr int start = -static_cast(rank); axis_type axes = index_sequence(start); - Plan plan(in, out, KOKKOS_FFT_BACKWARD, axes); + Plan plan(exec_space, in, out, KOKKOS_FFT_BACKWARD, axes); if(plan.is_transpose_needed()) { InViewType in_T; OutViewType out_T; - KokkosFFT::transpose(in, in_T, plan.map()); - KokkosFFT::transpose(out, out_T, plan.map()); + KokkosFFT::transpose(exec_space, in, in_T, plan.map()); + KokkosFFT::transpose(exec_space, out, out_T, plan.map()); - _ifft(plan, in_T, out_T, norm); + _ifft(exec_space, plan, in_T, out_T, norm); - KokkosFFT::transpose(out_T, out, plan.map_inv()); + KokkosFFT::transpose(exec_space, out_T, out, plan.map_inv()); } else { - _ifft(plan, in, out, norm); + _ifft(exec_space, plan, in, out, norm); } } - template - void ifftn(const InViewType& in, OutViewType& out, axis_type axes, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { + template + void ifftn(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, axis_type axes, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { static_assert(Kokkos::is_view::value, "KokkosFFT::ifftn: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, "KokkosFFT::ifftn: OutViewType is not a Kokkos::View."); - Plan plan(in, out, KOKKOS_FFT_BACKWARD, axes); + Plan plan(exec_space, in, out, KOKKOS_FFT_BACKWARD, axes); if(plan.is_transpose_needed()) { InViewType in_T; OutViewType out_T; - KokkosFFT::transpose(in, in_T, plan.map()); - KokkosFFT::transpose(out, out_T, plan.map()); + KokkosFFT::transpose(exec_space, in, in_T, plan.map()); + KokkosFFT::transpose(exec_space, out, out_T, plan.map()); - _ifft(plan, in_T, out_T, norm); + _ifft(exec_space, plan, in_T, out_T, norm); - KokkosFFT::transpose(out_T, out, plan.map_inv()); + KokkosFFT::transpose(exec_space, out_T, out, plan.map_inv()); } else { - _ifft(plan, in, out, norm); + _ifft(exec_space, plan, in, out, norm); } } - template - void rfftn(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { + template + void rfftn(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { static_assert(Kokkos::is_view::value, "KokkosFFT::rfftn: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -348,11 +356,11 @@ namespace KokkosFFT { static_assert(is_complex::value, "KokkosFFT::rfftn: OutViewType must be complex"); - fftn(in, out, norm); + fftn(exec_space, in, out, norm); } - template - void rfftn(const InViewType& in, OutViewType& out, axis_type axes, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { + template + void rfftn(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, axis_type axes, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { static_assert(Kokkos::is_view::value, "KokkosFFT::rfftn: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -366,11 +374,11 @@ namespace KokkosFFT { static_assert(is_complex::value, "KokkosFFT::rfftn: OutViewType must be complex"); - fftn(in, out, axes, norm); + fftn(exec_space, in, out, axes, norm); } - template - void irfftn(const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { + template + void irfftn(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { static_assert(Kokkos::is_view::value, "KokkosFFT::irfftn: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -384,11 +392,11 @@ namespace KokkosFFT { static_assert(std::is_floating_point::value, "KokkosFFT::irfftn: OutViewType must be real"); - ifftn(in, out, norm); + ifftn(exec_space, in, out, norm); } - template - void irfftn(const InViewType& in, OutViewType& out, axis_type axes, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { + template + void irfftn(const ExecutionSpace& exec_space, const InViewType& in, OutViewType& out, axis_type axes, KokkosFFT::Normalization norm=KokkosFFT::Normalization::BACKWARD) { static_assert(Kokkos::is_view::value, "KokkosFFT::irfftn: InViewType is not a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -402,7 +410,7 @@ namespace KokkosFFT { static_assert(std::is_floating_point::value, "KokkosFFT::irfftn: OutViewType must be real"); - ifftn(in, out, axes, norm); + ifftn(exec_space, in, out, axes, norm); } }; diff --git a/fft/unit_test/Test_Plans.cpp b/fft/unit_test/Test_Plans.cpp index ce759f45..c15bb358 100644 --- a/fft/unit_test/Test_Plans.cpp +++ b/fft/unit_test/Test_Plans.cpp @@ -16,24 +16,24 @@ void test_plan_1dfft_1dview() { ComplexView1DType x_cin("x_cin", n), x_cout("x_cout", n); // R2C plan - KokkosFFT::Plan plan_r2c(x, x_c); - KokkosFFT::Plan plan_r2c_axis_0(x, x_c, /*axis=*/0); - KokkosFFT::Plan plan_r2c_axes_0(x, x_c, /*axes=*/axes_type<1>({0})); + KokkosFFT::Plan plan_r2c(execution_space(), x, x_c); + KokkosFFT::Plan plan_r2c_axis_0(execution_space(), x, x_c, /*axis=*/0); + KokkosFFT::Plan plan_r2c_axes_0(execution_space(), x, x_c, /*axes=*/axes_type<1>({0})); // C2R plan - KokkosFFT::Plan plan_c2r(x_c, x); - KokkosFFT::Plan plan_c2r_axis0(x_c, x, /*axis=*/0); - KokkosFFT::Plan plan_c2r_axes0(x_c, x, /*axes=*/axes_type<1>({0})); + KokkosFFT::Plan plan_c2r(execution_space(), x_c, x); + KokkosFFT::Plan plan_c2r_axis0(execution_space(), x_c, x, /*axis=*/0); + KokkosFFT::Plan plan_c2r_axes0(execution_space(), x_c, x, /*axes=*/axes_type<1>({0})); // C2C plan - KokkosFFT::Plan plan_c2c_f(x_cin, x_cout, KOKKOS_FFT_FORWARD); - KokkosFFT::Plan plan_c2c_b(x_cin, x_cout, KOKKOS_FFT_BACKWARD); - KokkosFFT::Plan plan_c2c_f_axis0(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/0); - KokkosFFT::Plan plan_c2c_f_axes0(x_cin, x_cout, KOKKOS_FFT_BACKWARD, /*axes=*/axes_type<1>({0})); + KokkosFFT::Plan plan_c2c_f(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD); + KokkosFFT::Plan plan_c2c_b(execution_space(), x_cin, x_cout, KOKKOS_FFT_BACKWARD); + KokkosFFT::Plan plan_c2c_f_axis0(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/0); + KokkosFFT::Plan plan_c2c_f_axes0(execution_space(), x_cin, x_cout, KOKKOS_FFT_BACKWARD, /*axes=*/axes_type<1>({0})); EXPECT_THROW( { - KokkosFFT::Plan plan_c2c(x_cin, x_cout); + KokkosFFT::Plan plan_c2c(execution_space(), x_cin, x_cout); }, std::runtime_error ); @@ -65,38 +65,38 @@ void test_plan_2dfft_2dview() { ComplexView2DType x_cin("x_cin", n0, n1), x_cout("x_cout", n0, n1); // R2C plan - KokkosFFT::Plan plan_r2c(x, x_c_axis_1); - KokkosFFT::Plan plan_r2c_axes_0_1(x, x_c_axis_1, /*axes=*/axes_type<2>({0, 1})); - KokkosFFT::Plan plan_r2c_axes_1_0(x, x_c_axis_0, /*axes=*/axes_type<2>({1, 0})); + KokkosFFT::Plan plan_r2c(execution_space(), x, x_c_axis_1); + KokkosFFT::Plan plan_r2c_axes_0_1(execution_space(), x, x_c_axis_1, /*axes=*/axes_type<2>({0, 1})); + KokkosFFT::Plan plan_r2c_axes_1_0(execution_space(), x, x_c_axis_0, /*axes=*/axes_type<2>({1, 0})); // C2R plan - KokkosFFT::Plan plan_c2r(x_c_axis_1, x); - KokkosFFT::Plan plan_c2r_axes_0_1(x_c_axis_1, x, /*axes=*/axes_type<2>({0, 1})); - KokkosFFT::Plan plan_c2r_axes_1_0(x_c_axis_0, x, /*axes=*/axes_type<2>({1, 0})); + KokkosFFT::Plan plan_c2r(execution_space(), x_c_axis_1, x); + KokkosFFT::Plan plan_c2r_axes_0_1(execution_space(), x_c_axis_1, x, /*axes=*/axes_type<2>({0, 1})); + KokkosFFT::Plan plan_c2r_axes_1_0(execution_space(), x_c_axis_0, x, /*axes=*/axes_type<2>({1, 0})); // C2C plan - KokkosFFT::Plan plan_c2c_f(x_cin, x_cout, KOKKOS_FFT_FORWARD); - KokkosFFT::Plan plan_c2c_b(x_cin, x_cout, KOKKOS_FFT_BACKWARD); - KokkosFFT::Plan plan_c2c_f_axes_0_1(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({0, 1})); - KokkosFFT::Plan plan_c2c_f_axes_1_0(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({1, 0})); + KokkosFFT::Plan plan_c2c_f(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD); + KokkosFFT::Plan plan_c2c_b(execution_space(), x_cin, x_cout, KOKKOS_FFT_BACKWARD); + KokkosFFT::Plan plan_c2c_f_axes_0_1(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({0, 1})); + KokkosFFT::Plan plan_c2c_f_axes_1_0(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({1, 0})); EXPECT_THROW( { - KokkosFFT::Plan plan_c2c(x_cin, x_cout); + KokkosFFT::Plan plan_c2c(execution_space(), x_cin, x_cout); }, std::runtime_error ); EXPECT_THROW( { - KokkosFFT::Plan plan_c2c_axes_0_1(x_cin, x_cout, /*axes=*/axes_type<2>({0, 1})); + KokkosFFT::Plan plan_c2c_axes_0_1(execution_space(), x_cin, x_cout, /*axes=*/axes_type<2>({0, 1})); }, std::runtime_error ); EXPECT_THROW( { - KokkosFFT::Plan plan_c2c_axes_1_0(x_cin, x_cout, /*axes=*/axes_type<2>({1, 0})); + KokkosFFT::Plan plan_c2c_axes_1_0(execution_space(), x_cin, x_cout, /*axes=*/axes_type<2>({1, 0})); }, std::runtime_error ); @@ -129,39 +129,39 @@ void test_plan_3dfft_3dview() { ComplexView3DType x_cin("x_cin", n0, n1, n2), x_cout("x_cout", n0, n1, n2); // R2C plan - KokkosFFT::Plan plan_r2c_axes_0_1_2(x, x_c_axis_2, /*axes=*/axes_type<3>({0, 1, 2})); - KokkosFFT::Plan plan_r2c_axes_0_2_1(x, x_c_axis_1, /*axes=*/axes_type<3>({0, 2, 1})); - KokkosFFT::Plan plan_r2c_axes_1_0_2(x, x_c_axis_2, /*axes=*/axes_type<3>({1, 0, 2})); - KokkosFFT::Plan plan_r2c_axes_1_2_0(x, x_c_axis_0, /*axes=*/axes_type<3>({1, 2, 0})); - KokkosFFT::Plan plan_r2c_axes_2_0_1(x, x_c_axis_1, /*axes=*/axes_type<3>({2, 0, 1})); - KokkosFFT::Plan plan_r2c_axes_2_1_0(x, x_c_axis_0, /*axes=*/axes_type<3>({2, 1, 0})); + KokkosFFT::Plan plan_r2c_axes_0_1_2(execution_space(), x, x_c_axis_2, /*axes=*/axes_type<3>({0, 1, 2})); + KokkosFFT::Plan plan_r2c_axes_0_2_1(execution_space(), x, x_c_axis_1, /*axes=*/axes_type<3>({0, 2, 1})); + KokkosFFT::Plan plan_r2c_axes_1_0_2(execution_space(), x, x_c_axis_2, /*axes=*/axes_type<3>({1, 0, 2})); + KokkosFFT::Plan plan_r2c_axes_1_2_0(execution_space(), x, x_c_axis_0, /*axes=*/axes_type<3>({1, 2, 0})); + KokkosFFT::Plan plan_r2c_axes_2_0_1(execution_space(), x, x_c_axis_1, /*axes=*/axes_type<3>({2, 0, 1})); + KokkosFFT::Plan plan_r2c_axes_2_1_0(execution_space(), x, x_c_axis_0, /*axes=*/axes_type<3>({2, 1, 0})); // C2R plan - KokkosFFT::Plan plan_c2r_axes_0_1_2(x_c_axis_2, x, /*axes=*/axes_type<3>({0, 1, 2})); - KokkosFFT::Plan plan_c2r_axes_0_2_1(x_c_axis_1, x, /*axes=*/axes_type<3>({0, 2, 1})); - KokkosFFT::Plan plan_c2r_axes_1_0_2(x_c_axis_2, x, /*axes=*/axes_type<3>({1, 0, 2})); - KokkosFFT::Plan plan_c2r_axes_1_2_0(x_c_axis_0, x, /*axes=*/axes_type<3>({1, 2, 0})); - KokkosFFT::Plan plan_c2r_axes_2_0_1(x_c_axis_1, x, /*axes=*/axes_type<3>({2, 0, 1})); - KokkosFFT::Plan plan_c2r_axes_2_1_0(x_c_axis_0, x, /*axes=*/axes_type<3>({2, 1, 0})); + KokkosFFT::Plan plan_c2r_axes_0_1_2(execution_space(), x_c_axis_2, x, /*axes=*/axes_type<3>({0, 1, 2})); + KokkosFFT::Plan plan_c2r_axes_0_2_1(execution_space(), x_c_axis_1, x, /*axes=*/axes_type<3>({0, 2, 1})); + KokkosFFT::Plan plan_c2r_axes_1_0_2(execution_space(), x_c_axis_2, x, /*axes=*/axes_type<3>({1, 0, 2})); + KokkosFFT::Plan plan_c2r_axes_1_2_0(execution_space(), x_c_axis_0, x, /*axes=*/axes_type<3>({1, 2, 0})); + KokkosFFT::Plan plan_c2r_axes_2_0_1(execution_space(), x_c_axis_1, x, /*axes=*/axes_type<3>({2, 0, 1})); + KokkosFFT::Plan plan_c2r_axes_2_1_0(execution_space(), x_c_axis_0, x, /*axes=*/axes_type<3>({2, 1, 0})); // C2C plan - KokkosFFT::Plan plan_c2c_f_axes_0_1_2(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({0, 1, 2})); - KokkosFFT::Plan plan_c2c_f_axes_0_2_1(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({0, 2, 1})); - KokkosFFT::Plan plan_c2c_f_axes_1_0_2(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({1, 0, 2})); - KokkosFFT::Plan plan_c2c_f_axes_1_2_0(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({1, 2, 0})); - KokkosFFT::Plan plan_c2c_f_axes_2_0_1(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({2, 0, 1})); - KokkosFFT::Plan plan_c2c_f_axes_2_1_0(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({2, 1, 0})); + KokkosFFT::Plan plan_c2c_f_axes_0_1_2(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({0, 1, 2})); + KokkosFFT::Plan plan_c2c_f_axes_0_2_1(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({0, 2, 1})); + KokkosFFT::Plan plan_c2c_f_axes_1_0_2(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({1, 0, 2})); + KokkosFFT::Plan plan_c2c_f_axes_1_2_0(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({1, 2, 0})); + KokkosFFT::Plan plan_c2c_f_axes_2_0_1(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({2, 0, 1})); + KokkosFFT::Plan plan_c2c_f_axes_2_1_0(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<3>({2, 1, 0})); EXPECT_THROW( { - KokkosFFT::Plan plan_c2c(x_cin, x_cout); + KokkosFFT::Plan plan_c2c(execution_space(), x_cin, x_cout); }, std::runtime_error ); EXPECT_THROW( { - KokkosFFT::Plan plan_c2c_axes_0_1_2(x_cin, x_cout, /*axes=*/axes_type<3>({0, 1, 2})); + KokkosFFT::Plan plan_c2c_axes_0_1_2(execution_space(), x_cin, x_cout, /*axes=*/axes_type<3>({0, 1, 2})); }, std::runtime_error ); @@ -193,31 +193,31 @@ void test_plan_1dfft_2dview() { ComplexView2DType x_cin("x_cin", n0, n1), x_cout("x_cout", n0, n1); // R2C plan - KokkosFFT::Plan plan_r2c_axis_0(x, x_c_axis_0, /*axis=*/0); - KokkosFFT::Plan plan_r2c_axis_1(x, x_c_axis_1, /*axis=*/1); - KokkosFFT::Plan plan_r2c_axis_minus1(x, x_c_axis_1, /*axis=*/-1); + KokkosFFT::Plan plan_r2c_axis_0(execution_space(), x, x_c_axis_0, /*axis=*/0); + KokkosFFT::Plan plan_r2c_axis_1(execution_space(), x, x_c_axis_1, /*axis=*/1); + KokkosFFT::Plan plan_r2c_axis_minus1(execution_space(), x, x_c_axis_1, /*axis=*/-1); // C2R plan - KokkosFFT::Plan plan_c2r_axis_0(x_c_axis_0, x, /*axis=*/0); - KokkosFFT::Plan plan_c2r_axis_1(x_c_axis_1, x, /*axis=*/1); - KokkosFFT::Plan plan_c2r_axis_minus1(x_c_axis_1, x, /*axis=*/-1); + KokkosFFT::Plan plan_c2r_axis_0(execution_space(), x_c_axis_0, x, /*axis=*/0); + KokkosFFT::Plan plan_c2r_axis_1(execution_space(), x_c_axis_1, x, /*axis=*/1); + KokkosFFT::Plan plan_c2r_axis_minus1(execution_space(), x_c_axis_1, x, /*axis=*/-1); // C2C plan - KokkosFFT::Plan plan_c2c_f_axis_0(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/0); - KokkosFFT::Plan plan_c2c_f_axis_1(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/1); + KokkosFFT::Plan plan_c2c_f_axis_0(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/0); + KokkosFFT::Plan plan_c2c_f_axis_1(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/1); // [TO DO] Fix this, this can be instanized with explicit Plan(InViewType& in, OutViewType& out, FFTDirectionType direction) // Because FFTDirectionType is int for most libraries //EXPECT_THROW( // { - // KokkosFFT::Plan plan_c2c_axis_0(x_cin, x_cout, /*axis=*/0); + // KokkosFFT::Plan plan_c2c_axis_0(execution_space(), x_cin, x_cout, /*axis=*/0); // }, // std::runtime_error //); //EXPECT_THROW( // { - // KokkosFFT::Plan plan_c2c_axis_1(x_cin, x_cout, /*axis=*/1); + // KokkosFFT::Plan plan_c2c_axis_1(execution_space(), x_cin, x_cout, /*axis=*/1); // }, // std::runtime_error //); @@ -250,42 +250,42 @@ void test_plan_1dfft_3dview() { ComplexView3DType x_cin("x_cin", n0, n1, n2), x_cout("x_cout", n0, n1, n2); // R2C plan - KokkosFFT::Plan plan_r2c_axis_0(x, x_c_axis_0, /*axis=*/0); - KokkosFFT::Plan plan_r2c_axis_1(x, x_c_axis_1, /*axis=*/1); - KokkosFFT::Plan plan_r2c_axis_2(x, x_c_axis_2, /*axis=*/2); + KokkosFFT::Plan plan_r2c_axis_0(execution_space(), x, x_c_axis_0, /*axis=*/0); + KokkosFFT::Plan plan_r2c_axis_1(execution_space(), x, x_c_axis_1, /*axis=*/1); + KokkosFFT::Plan plan_r2c_axis_2(execution_space(), x, x_c_axis_2, /*axis=*/2); // C2R plan - KokkosFFT::Plan plan_c2r_axis_0(x_c_axis_0, x, /*axis=*/0); - KokkosFFT::Plan plan_c2r_axis_1(x_c_axis_1, x, /*axis=*/1); - KokkosFFT::Plan plan_c2r_axis_2(x_c_axis_2, x, /*axis=*/2); + KokkosFFT::Plan plan_c2r_axis_0(execution_space(), x_c_axis_0, x, /*axis=*/0); + KokkosFFT::Plan plan_c2r_axis_1(execution_space(), x_c_axis_1, x, /*axis=*/1); + KokkosFFT::Plan plan_c2r_axis_2(execution_space(), x_c_axis_2, x, /*axis=*/2); // C2C plan - KokkosFFT::Plan plan_c2c_f_axis_0(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/0); - KokkosFFT::Plan plan_c2c_f_axis_1(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/1); - KokkosFFT::Plan plan_c2c_f_axis_2(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/2); - KokkosFFT::Plan plan_c2c_b_axis_0(x_cin, x_cout, KOKKOS_FFT_BACKWARD, /*axis=*/0); - KokkosFFT::Plan plan_c2c_b_axis_1(x_cin, x_cout, KOKKOS_FFT_BACKWARD, /*axis=*/1); - KokkosFFT::Plan plan_c2c_b_axis_2(x_cin, x_cout, KOKKOS_FFT_BACKWARD, /*axis=*/2); + KokkosFFT::Plan plan_c2c_f_axis_0(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/0); + KokkosFFT::Plan plan_c2c_f_axis_1(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/1); + KokkosFFT::Plan plan_c2c_f_axis_2(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axis=*/2); + KokkosFFT::Plan plan_c2c_b_axis_0(execution_space(), x_cin, x_cout, KOKKOS_FFT_BACKWARD, /*axis=*/0); + KokkosFFT::Plan plan_c2c_b_axis_1(execution_space(), x_cin, x_cout, KOKKOS_FFT_BACKWARD, /*axis=*/1); + KokkosFFT::Plan plan_c2c_b_axis_2(execution_space(), x_cin, x_cout, KOKKOS_FFT_BACKWARD, /*axis=*/2); // [TO DO] Fix this, this can be instanized with explicit Plan(InViewType& in, OutViewType& out, FFTDirectionType direction) // Because FFTDirectionType is int for most libraries //EXPECT_THROW( // { - // KokkosFFT::Plan plan_c2c_axis_0(x_cin, x_cout, /*axis=*/0); + // KokkosFFT::Plan plan_c2c_axis_0(execution_space(), x_cin, x_cout, /*axis=*/0); // }, // std::runtime_error //); //EXPECT_THROW( // { - // KokkosFFT::Plan plan_c2c_axis_1(x_cin, x_cout, /*axis=*/1); + // KokkosFFT::Plan plan_c2c_axis_1(execution_space(), x_cin, x_cout, /*axis=*/1); // }, // std::runtime_error //); //EXPECT_THROW( // { - // KokkosFFT::Plan plan_c2c_axis_1(x_cin, x_cout, /*axis=*/2); + // KokkosFFT::Plan plan_c2c_axis_1(execution_space(), x_cin, x_cout, /*axis=*/2); // }, // std::runtime_error //); @@ -318,32 +318,32 @@ void test_plan_2dfft_3dview() { ComplexView3DType x_cin("x_cin", n0, n1, n2), x_cout("x_cout", n0, n1, n2); // R2C plan - KokkosFFT::Plan plan_r2c_axes_0_1(x, x_c_axis_1, /*axes=*/axes_type<2>({0, 1})); - KokkosFFT::Plan plan_r2c_axes_0_2(x, x_c_axis_2, /*axes=*/axes_type<2>({0, 2})); - KokkosFFT::Plan plan_r2c_axes_1_0(x, x_c_axis_0, /*axes=*/axes_type<2>({1, 0})); - KokkosFFT::Plan plan_r2c_axes_1_2(x, x_c_axis_2, /*axes=*/axes_type<2>({1, 2})); - KokkosFFT::Plan plan_r2c_axes_2_0(x, x_c_axis_0, /*axes=*/axes_type<2>({2, 0})); - KokkosFFT::Plan plan_r2c_axes_2_1(x, x_c_axis_1, /*axes=*/axes_type<2>({2, 1})); + KokkosFFT::Plan plan_r2c_axes_0_1(execution_space(), x, x_c_axis_1, /*axes=*/axes_type<2>({0, 1})); + KokkosFFT::Plan plan_r2c_axes_0_2(execution_space(), x, x_c_axis_2, /*axes=*/axes_type<2>({0, 2})); + KokkosFFT::Plan plan_r2c_axes_1_0(execution_space(), x, x_c_axis_0, /*axes=*/axes_type<2>({1, 0})); + KokkosFFT::Plan plan_r2c_axes_1_2(execution_space(), x, x_c_axis_2, /*axes=*/axes_type<2>({1, 2})); + KokkosFFT::Plan plan_r2c_axes_2_0(execution_space(), x, x_c_axis_0, /*axes=*/axes_type<2>({2, 0})); + KokkosFFT::Plan plan_r2c_axes_2_1(execution_space(), x, x_c_axis_1, /*axes=*/axes_type<2>({2, 1})); // C2R plan - KokkosFFT::Plan plan_c2r_axes_0_1(x_c_axis_1, x, /*axes=*/axes_type<2>({0, 1})); - KokkosFFT::Plan plan_c2r_axes_0_2(x_c_axis_2, x, /*axes=*/axes_type<2>({0, 2})); - KokkosFFT::Plan plan_c2r_axes_1_0(x_c_axis_0, x, /*axes=*/axes_type<2>({1, 0})); - KokkosFFT::Plan plan_c2r_axes_1_2(x_c_axis_2, x, /*axes=*/axes_type<2>({1, 2})); - KokkosFFT::Plan plan_c2r_axes_2_0(x_c_axis_0, x, /*axes=*/axes_type<2>({2, 0})); - KokkosFFT::Plan plan_c2r_axes_2_1(x_c_axis_1, x, /*axes=*/axes_type<2>({2, 1})); + KokkosFFT::Plan plan_c2r_axes_0_1(execution_space(), x_c_axis_1, x, /*axes=*/axes_type<2>({0, 1})); + KokkosFFT::Plan plan_c2r_axes_0_2(execution_space(), x_c_axis_2, x, /*axes=*/axes_type<2>({0, 2})); + KokkosFFT::Plan plan_c2r_axes_1_0(execution_space(), x_c_axis_0, x, /*axes=*/axes_type<2>({1, 0})); + KokkosFFT::Plan plan_c2r_axes_1_2(execution_space(), x_c_axis_2, x, /*axes=*/axes_type<2>({1, 2})); + KokkosFFT::Plan plan_c2r_axes_2_0(execution_space(), x_c_axis_0, x, /*axes=*/axes_type<2>({2, 0})); + KokkosFFT::Plan plan_c2r_axes_2_1(execution_space(), x_c_axis_1, x, /*axes=*/axes_type<2>({2, 1})); // C2C plan - KokkosFFT::Plan plan_c2c_f_axes_0_1(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({0, 1})); - KokkosFFT::Plan plan_c2c_f_axes_0_2(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({0, 2})); - KokkosFFT::Plan plan_c2c_f_axes_1_0(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({1, 0})); - KokkosFFT::Plan plan_c2c_f_axes_1_2(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({1, 2})); - KokkosFFT::Plan plan_c2c_f_axes_2_0(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({2, 0})); - KokkosFFT::Plan plan_c2c_f_axes_2_1(x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({2, 1})); + KokkosFFT::Plan plan_c2c_f_axes_0_1(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({0, 1})); + KokkosFFT::Plan plan_c2c_f_axes_0_2(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({0, 2})); + KokkosFFT::Plan plan_c2c_f_axes_1_0(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({1, 0})); + KokkosFFT::Plan plan_c2c_f_axes_1_2(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({1, 2})); + KokkosFFT::Plan plan_c2c_f_axes_2_0(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({2, 0})); + KokkosFFT::Plan plan_c2c_f_axes_2_1(execution_space(), x_cin, x_cout, KOKKOS_FFT_FORWARD, /*axes=*/axes_type<2>({2, 1})); EXPECT_THROW( { - KokkosFFT::Plan plan_c2c_axes_0_1(x_cin, x_cout, /*axes=*/axes_type<2>({0, 1})); + KokkosFFT::Plan plan_c2c_axes_0_1(execution_space(), x_cin, x_cout, /*axes=*/axes_type<2>({0, 1})); }, std::runtime_error ); diff --git a/fft/unit_test/Test_Transform.cpp b/fft/unit_test/Test_Transform.cpp index a52f90ea..56e2d20b 100644 --- a/fft/unit_test/Test_Transform.cpp +++ b/fft/unit_test/Test_Transform.cpp @@ -103,11 +103,11 @@ void test_fft1_identity(T atol=1.0e-12) { Kokkos::fence(); - KokkosFFT::fft(a, out); - KokkosFFT::ifft(out, _a); + KokkosFFT::fft(execution_space(), a, out); + KokkosFFT::ifft(execution_space(), out, _a); - KokkosFFT::rfft(ar, outr); - KokkosFFT::irfft(outr, _ar); + KokkosFFT::rfft(execution_space(), ar, outr); + KokkosFFT::irfft(execution_space(), outr, _ar); EXPECT_TRUE( allclose(_a, a_ref, 1.e-5, atol) ); EXPECT_TRUE( allclose(_ar, ar_ref, 1.e-5, atol) ); @@ -127,10 +127,10 @@ void test_fft1_1dfft_1dview() { Kokkos::fence(); - KokkosFFT::fft(x, out); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::fft(x, out_b, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::fft(x, out_o, KokkosFFT::Normalization::ORTHO); - KokkosFFT::fft(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::fft(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::fft(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::fft(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::fft(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); fft1(x, ref); multiply(out_o, sqrt(static_cast(len))); @@ -156,10 +156,10 @@ void test_fft1_1difft_1dview() { Kokkos::fence(); - KokkosFFT::ifft(x, out); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::ifft(x, out_b, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::ifft(x, out_o, KokkosFFT::Normalization::ORTHO); - KokkosFFT::ifft(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::ifft(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::ifft(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::ifft(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::ifft(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); ifft1(x, ref); multiply(out_o, sqrt(static_cast(len))); @@ -206,15 +206,15 @@ void test_fft1_1dfft_2dview(T atol=1.e-12) { fft1(sub_x, sub_ref); } - KokkosFFT::fft(x, out_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::fft(execution_space(), x, out_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); EXPECT_TRUE( allclose(out_axis0, ref_out_axis0, 1.e-5, atol) ); - KokkosFFT::ifft(out_axis0, x_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::ifft(execution_space(), out_axis0, x_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); EXPECT_TRUE( allclose(x_axis0, ref_x, 1.e-5, atol) ); // Simple identity tests for r2c and c2r transforms - KokkosFFT::rfft(xr, outr_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); - KokkosFFT::irfft(outr_axis0, xr_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::rfft(execution_space(), xr, outr_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::irfft(execution_space(), outr_axis0, xr_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); EXPECT_TRUE( allclose(xr_axis0, ref_xr, 1.e-5, atol) ); @@ -230,15 +230,15 @@ void test_fft1_1dfft_2dview(T atol=1.e-12) { fft1(sub_x, sub_ref); } - KokkosFFT::fft(x, out_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::fft(execution_space(), x, out_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); EXPECT_TRUE( allclose(out_axis1, ref_out_axis1, 1.e-5, atol) ); - KokkosFFT::ifft(out_axis1, x_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::ifft(execution_space(), out_axis1, x_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); EXPECT_TRUE( allclose(x_axis1, ref_x, 1.e-5, atol) ); // Simple identity tests for r2c and c2r transforms - KokkosFFT::rfft(xr, outr_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::irfft(outr_axis1, xr_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::rfft(execution_space(), xr, outr_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::irfft(execution_space(), outr_axis1, xr_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); EXPECT_TRUE( allclose(xr_axis1, ref_xr, 1.e-5, atol) ); } @@ -280,15 +280,15 @@ void test_fft1_1dfft_3dview(T atol=1.e-12) { } } - KokkosFFT::fft(x, out_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::fft(execution_space(), x, out_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); EXPECT_TRUE( allclose(out_axis0, ref_out_axis0, 1.e-5, atol) ); - KokkosFFT::ifft(out_axis0, x_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::ifft(execution_space(), out_axis0, x_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); EXPECT_TRUE( allclose(x_axis0, ref_x, 1.e-5, atol) ); // Simple identity tests for r2c and c2r transforms - KokkosFFT::rfft(xr, outr_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); - KokkosFFT::irfft(outr_axis0, xr_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::rfft(execution_space(), xr, outr_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::irfft(execution_space(), outr_axis0, xr_axis0, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); EXPECT_TRUE( allclose(xr_axis0, ref_xr, 1.e-5, atol) ); @@ -306,15 +306,15 @@ void test_fft1_1dfft_3dview(T atol=1.e-12) { } } - KokkosFFT::fft(x, out_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::fft(execution_space(), x, out_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); EXPECT_TRUE( allclose(out_axis1, ref_out_axis1, 1.e-5, atol) ); - KokkosFFT::ifft(out_axis1, x_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::ifft(execution_space(), out_axis1, x_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); EXPECT_TRUE( allclose(x_axis1, ref_x, 1.e-5, atol) ); // Simple identity tests for r2c and c2r transforms - KokkosFFT::rfft(xr, outr_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::irfft(outr_axis1, xr_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::rfft(execution_space(), xr, outr_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::irfft(execution_space(), outr_axis1, xr_axis1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); EXPECT_TRUE( allclose(xr_axis1, ref_xr, 1.e-5, atol) ); @@ -332,15 +332,15 @@ void test_fft1_1dfft_3dview(T atol=1.e-12) { } } - KokkosFFT::fft(x, out_axis2, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); + KokkosFFT::fft(execution_space(), x, out_axis2, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); EXPECT_TRUE( allclose(out_axis2, ref_out_axis2, 1.e-5, atol) ); - KokkosFFT::ifft(out_axis2, x_axis2, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); + KokkosFFT::ifft(execution_space(), out_axis2, x_axis2, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); EXPECT_TRUE( allclose(x_axis2, ref_x, 1.e-5, atol) ); // Simple identity tests for r2c and c2r transforms - KokkosFFT::rfft(xr, outr_axis2, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); - KokkosFFT::irfft(outr_axis2, xr_axis2, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); + KokkosFFT::rfft(execution_space(), xr, outr_axis2, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); + KokkosFFT::irfft(execution_space(), outr_axis2, xr_axis2, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); EXPECT_TRUE( allclose(xr_axis2, ref_xr, 1.e-5, atol) ); } @@ -447,13 +447,13 @@ void test_fft2_2dfft_2dview() { Kokkos::fence(); // np.fft2 is identical to np.fft(np.fft(x, axis=1), axis=0) - KokkosFFT::fft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::fft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::fft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::fft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); - KokkosFFT::fft2(x, out); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::fft2(x, out_b, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::fft2(x, out_o, KokkosFFT::Normalization::ORTHO); - KokkosFFT::fft2(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::fft2(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::fft2(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::fft2(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::fft2(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, sqrt(static_cast(n0 * n1))); multiply(out_f, static_cast(n0 * n1)); @@ -480,13 +480,13 @@ void test_fft2_2difft_2dview() { Kokkos::fence(); // np.ifft2 is identical to np.ifft(np.ifft(x, axis=1), axis=0) - KokkosFFT::ifft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::ifft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::ifft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::ifft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); - KokkosFFT::ifft2(x, out); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::ifft2(x, out_b, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::ifft2(x, out_o, KokkosFFT::Normalization::ORTHO); - KokkosFFT::ifft2(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::ifft2(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::ifft2(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::ifft2(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::ifft2(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, 1.0/sqrt(static_cast(n0 * n1))); multiply(out_f, 1.0/static_cast(n0 * n1)); @@ -513,20 +513,20 @@ void test_fft2_2drfft_2dview() { Kokkos::fence(); // np.rfft2 is identical to np.fft(np.rfft(x, axis=1), axis=0) - KokkosFFT::rfft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::fft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::rfft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::fft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfft2(x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::rfft2(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfft2(x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::rfft2(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfft2(x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::rfft2(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfft2(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::rfft2(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, sqrt(static_cast(n0 * n1))); multiply(out_f, static_cast(n0 * n1)); @@ -554,20 +554,20 @@ void test_fft2_2dirfft_2dview() { Kokkos::deep_copy(x_ref, x); // np.irfft2 is identical to np.irfft(np.ifft(x, axis=0), axis=1) - KokkosFFT::ifft(x, out1, KokkosFFT::Normalization::BACKWARD, 0); - KokkosFFT::irfft(out1, out2, KokkosFFT::Normalization::BACKWARD, 1); + KokkosFFT::ifft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, 0); + KokkosFFT::irfft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, 1); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfft2(x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::irfft2(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfft2(x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::irfft2(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfft2(x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::irfft2(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfft2(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::irfft2(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, 1.0/sqrt(static_cast(n0 * n1))); multiply(out_f, 1.0/static_cast(n0 * n1)); @@ -663,13 +663,13 @@ void test_fftn_2dfft_2dview() { Kokkos::fence(); // np.fftn for 2D array is identical to np.fft(np.fft(x, axis=1), axis=0) - KokkosFFT::fft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::fft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::fft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::fft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); - KokkosFFT::fftn(x, out); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::fftn(x, out_b, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::fftn(x, out_o, KokkosFFT::Normalization::ORTHO); - KokkosFFT::fftn(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::fftn(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::fftn(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::fftn(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::fftn(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, sqrt(static_cast(n0 * n1))); multiply(out_f, static_cast(n0 * n1)); @@ -683,10 +683,10 @@ void test_fftn_2dfft_2dview() { // np.fftn for 2D array is identical to np.fft(np.fft(x, axis=1), axis=0) using axes_type = KokkosFFT::axis_type<2>; - KokkosFFT::fftn(x, out, axes_type{-2, -1}); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::fftn(x, out_b, axes_type{-2, -1}, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::fftn(x, out_o, axes_type{-2, -1}, KokkosFFT::Normalization::ORTHO); - KokkosFFT::fftn(x, out_f, axes_type{-2, -1}, KokkosFFT::Normalization::FORWARD); + KokkosFFT::fftn(execution_space(), x, out, axes_type{-2, -1}); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::fftn(execution_space(), x, out_b, axes_type{-2, -1}, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::fftn(execution_space(), x, out_o, axes_type{-2, -1}, KokkosFFT::Normalization::ORTHO); + KokkosFFT::fftn(execution_space(), x, out_f, axes_type{-2, -1}, KokkosFFT::Normalization::FORWARD); multiply(out_o, sqrt(static_cast(n0 * n1))); multiply(out_f, static_cast(n0 * n1)); @@ -713,14 +713,14 @@ void test_fftn_3dfft_3dview(T atol=1.0e-6) { Kokkos::fence(); // np.fftn for 3D array is identical to np.fft(np.fft(np.fft(x, axis=2), axis=1), axis=0) - KokkosFFT::fft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); - KokkosFFT::fft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::fft(out2, out3, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::fft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); + KokkosFFT::fft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::fft(execution_space(), out2, out3, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); - KokkosFFT::fftn(x, out); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::fftn(x, out_b, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::fftn(x, out_o, KokkosFFT::Normalization::ORTHO); - KokkosFFT::fftn(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::fftn(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::fftn(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::fftn(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::fftn(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, sqrt(static_cast(n0 * n1 * n2))); multiply(out_f, static_cast(n0 * n1 * n2)); @@ -734,10 +734,10 @@ void test_fftn_3dfft_3dview(T atol=1.0e-6) { // np.fftn for 3D array is identical to np.fft(np.fft(np.fft(x, axis=2), axis=1), axis=0) using axes_type = KokkosFFT::axis_type<3>; - KokkosFFT::fftn(x, out, axes_type{-3, -2, -1}); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::fftn(x, out_b, axes_type{-3, -2, -1}, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::fftn(x, out_o, axes_type{-3, -2, -1}, KokkosFFT::Normalization::ORTHO); - KokkosFFT::fftn(x, out_f, axes_type{-3, -2, -1}, KokkosFFT::Normalization::FORWARD); + KokkosFFT::fftn(execution_space(), x, out, axes_type{-3, -2, -1}); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::fftn(execution_space(), x, out_b, axes_type{-3, -2, -1}, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::fftn(execution_space(), x, out_o, axes_type{-3, -2, -1}, KokkosFFT::Normalization::ORTHO); + KokkosFFT::fftn(execution_space(), x, out_f, axes_type{-3, -2, -1}, KokkosFFT::Normalization::FORWARD); multiply(out_o, sqrt(static_cast(n0 * n1 * n2))); multiply(out_f, static_cast(n0 * n1 * n2)); @@ -764,13 +764,13 @@ void test_ifftn_2dfft_2dview() { Kokkos::fence(); // np.ifftn for 2D array is identical to np.ifft(np.ifft(x, axis=1), axis=0) - KokkosFFT::ifft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::ifft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::ifft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::ifft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); - KokkosFFT::ifftn(x, out); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::ifftn(x, out_b, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::ifftn(x, out_o, KokkosFFT::Normalization::ORTHO); - KokkosFFT::ifftn(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::ifftn(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::ifftn(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::ifftn(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::ifftn(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, 1.0/sqrt(static_cast(n0 * n1))); multiply(out_f, 1.0/static_cast(n0 * n1)); @@ -784,10 +784,10 @@ void test_ifftn_2dfft_2dview() { // np.fftn for 2D array is identical to np.fft(np.fft(x, axis=1), axis=0) using axes_type = KokkosFFT::axis_type<2>; - KokkosFFT::ifftn(x, out, axes_type{-2, -1}); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::ifftn(x, out_b, axes_type{-2, -1}, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::ifftn(x, out_o, axes_type{-2, -1}, KokkosFFT::Normalization::ORTHO); - KokkosFFT::ifftn(x, out_f, axes_type{-2, -1}, KokkosFFT::Normalization::FORWARD); + KokkosFFT::ifftn(execution_space(), x, out, axes_type{-2, -1}); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::ifftn(execution_space(), x, out_b, axes_type{-2, -1}, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::ifftn(execution_space(), x, out_o, axes_type{-2, -1}, KokkosFFT::Normalization::ORTHO); + KokkosFFT::ifftn(execution_space(), x, out_f, axes_type{-2, -1}, KokkosFFT::Normalization::FORWARD); multiply(out_o, 1.0/sqrt(static_cast(n0 * n1))); multiply(out_f, 1.0/static_cast(n0 * n1)); @@ -814,14 +814,14 @@ void test_ifftn_3dfft_3dview() { Kokkos::fence(); // np.ifftn for 3D array is identical to np.ifft(np.ifft(np.ifft(x, axis=2), axis=1), axis=0) - KokkosFFT::ifft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); - KokkosFFT::ifft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::ifft(out2, out3, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::ifft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); + KokkosFFT::ifft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::ifft(execution_space(), out2, out3, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); - KokkosFFT::ifftn(x, out); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::ifftn(x, out_b, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::ifftn(x, out_o, KokkosFFT::Normalization::ORTHO); - KokkosFFT::ifftn(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::ifftn(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::ifftn(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::ifftn(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::ifftn(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, 1.0/sqrt(static_cast(n0 * n1 * n2))); multiply(out_f, 1.0/static_cast(n0 * n1 * n2)); @@ -835,10 +835,10 @@ void test_ifftn_3dfft_3dview() { // np.ifftn for 3D array is identical to np.ifft(np.ifft(np.ifft(x, axis=2), axis=1), axis=0) using axes_type = KokkosFFT::axis_type<3>; - KokkosFFT::ifftn(x, out, axes_type{-3, -2, -1}); // default: KokkosFFT::Normalization::BACKWARD - KokkosFFT::ifftn(x, out_b, axes_type{-3, -2, -1}, KokkosFFT::Normalization::BACKWARD); - KokkosFFT::ifftn(x, out_o, axes_type{-3, -2, -1}, KokkosFFT::Normalization::ORTHO); - KokkosFFT::ifftn(x, out_f, axes_type{-3, -2, -1}, KokkosFFT::Normalization::FORWARD); + KokkosFFT::ifftn(execution_space(), x, out, axes_type{-3, -2, -1}); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::ifftn(execution_space(), x, out_b, axes_type{-3, -2, -1}, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::ifftn(execution_space(), x, out_o, axes_type{-3, -2, -1}, KokkosFFT::Normalization::ORTHO); + KokkosFFT::ifftn(execution_space(), x, out_f, axes_type{-3, -2, -1}, KokkosFFT::Normalization::FORWARD); multiply(out_o, 1.0/sqrt(static_cast(n0 * n1 * n2))); multiply(out_f, 1.0/static_cast(n0 * n1 * n2)); @@ -865,20 +865,20 @@ void test_rfftn_2dfft_2dview() { Kokkos::fence(); // np.rfftn for 2D array is identical to np.fft(np.rfft(x, axis=1), axis=0) - KokkosFFT::rfft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::fft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::rfft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::fft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::rfftn(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::rfftn(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::rfftn(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::rfftn(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, sqrt(static_cast(n0 * n1))); multiply(out_f, static_cast(n0 * n1)); @@ -893,16 +893,16 @@ void test_rfftn_2dfft_2dview() { using axes_type = KokkosFFT::axis_type<2>; Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out, axes_type{-2, -1}); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::rfftn(execution_space(), x, out, axes_type{-2, -1}); // default: KokkosFFT::Normalization::BACKWARD Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_b, axes_type{-2, -1}, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::rfftn(execution_space(), x, out_b, axes_type{-2, -1}, KokkosFFT::Normalization::BACKWARD); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_o, axes_type{-2, -1}, KokkosFFT::Normalization::ORTHO); + KokkosFFT::rfftn(execution_space(), x, out_o, axes_type{-2, -1}, KokkosFFT::Normalization::ORTHO); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_f, axes_type{-2, -1}, KokkosFFT::Normalization::FORWARD); + KokkosFFT::rfftn(execution_space(), x, out_f, axes_type{-2, -1}, KokkosFFT::Normalization::FORWARD); multiply(out_o, sqrt(static_cast(n0 * n1))); multiply(out_f, static_cast(n0 * n1)); @@ -929,21 +929,21 @@ void test_rfftn_3dfft_3dview() { Kokkos::fence(); // np.rfftn for 3D array is identical to np.fft(np.fft(np.rfft(x, axis=2), axis=1), axis=0) - KokkosFFT::rfft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); - KokkosFFT::fft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::fft(out2, out3, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::rfft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); + KokkosFFT::fft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::fft(execution_space(), out2, out3, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::rfftn(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::rfftn(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::rfftn(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::rfftn(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, sqrt(static_cast(n0 * n1 * n2))); multiply(out_f, static_cast(n0 * n1 * n2)); @@ -958,16 +958,16 @@ void test_rfftn_3dfft_3dview() { using axes_type = KokkosFFT::axis_type<3>; Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out, axes_type{-3, -2, -1}); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::rfftn(execution_space(), x, out, axes_type{-3, -2, -1}); // default: KokkosFFT::Normalization::BACKWARD Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_b, axes_type{-3, -2, -1}, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::rfftn(execution_space(), x, out_b, axes_type{-3, -2, -1}, KokkosFFT::Normalization::BACKWARD); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_o, axes_type{-3, -2, -1}, KokkosFFT::Normalization::ORTHO); + KokkosFFT::rfftn(execution_space(), x, out_o, axes_type{-3, -2, -1}, KokkosFFT::Normalization::ORTHO); Kokkos::deep_copy(x, x_ref); - KokkosFFT::rfftn(x, out_f, axes_type{-3, -2, -1}, KokkosFFT::Normalization::FORWARD); + KokkosFFT::rfftn(execution_space(), x, out_f, axes_type{-3, -2, -1}, KokkosFFT::Normalization::FORWARD); multiply(out_o, sqrt(static_cast(n0 * n1 * n2))); multiply(out_f, static_cast(n0 * n1 * n2)); @@ -995,20 +995,20 @@ void test_irfftn_2dfft_2dview() { Kokkos::deep_copy(x_ref, x); // np.irfftn for 2D array is identical to np.irfft(np.ifft(x, axis=0), axis=1) - KokkosFFT::ifft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); - KokkosFFT::irfft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::ifft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::irfft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::irfftn(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::irfftn(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::irfftn(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::irfftn(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, 1.0/sqrt(static_cast(n0 * n1))); multiply(out_f, 1.0/static_cast(n0 * n1)); @@ -1023,16 +1023,16 @@ void test_irfftn_2dfft_2dview() { using axes_type = KokkosFFT::axis_type<2>; Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out, axes_type{-2, -1}); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::irfftn(execution_space(), x, out, axes_type{-2, -1}); // default: KokkosFFT::Normalization::BACKWARD Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_b, axes_type{-2, -1}, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::irfftn(execution_space(), x, out_b, axes_type{-2, -1}, KokkosFFT::Normalization::BACKWARD); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_o, axes_type{-2, -1}, KokkosFFT::Normalization::ORTHO); + KokkosFFT::irfftn(execution_space(), x, out_o, axes_type{-2, -1}, KokkosFFT::Normalization::ORTHO); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_f, axes_type{-2, -1}, KokkosFFT::Normalization::FORWARD); + KokkosFFT::irfftn(execution_space(), x, out_f, axes_type{-2, -1}, KokkosFFT::Normalization::FORWARD); multiply(out_o, 1.0/sqrt(static_cast(n0 * n1))); multiply(out_f, 1.0/static_cast(n0 * n1)); @@ -1060,21 +1060,21 @@ void test_irfftn_3dfft_3dview() { Kokkos::deep_copy(x_ref, x); // np.irfftn for 3D array is identical to np.irfft(np.ifft(np.ifft(x, axis=0), axis=1), axis=2) - KokkosFFT::ifft(x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); - KokkosFFT::ifft(out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); - KokkosFFT::irfft(out2, out3, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); + KokkosFFT::ifft(execution_space(), x, out1, KokkosFFT::Normalization::BACKWARD, /*axis=*/0); + KokkosFFT::ifft(execution_space(), out1, out2, KokkosFFT::Normalization::BACKWARD, /*axis=*/1); + KokkosFFT::irfft(execution_space(), out2, out3, KokkosFFT::Normalization::BACKWARD, /*axis=*/2); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::irfftn(execution_space(), x, out); // default: KokkosFFT::Normalization::BACKWARD Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_b, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::irfftn(execution_space(), x, out_b, KokkosFFT::Normalization::BACKWARD); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_o, KokkosFFT::Normalization::ORTHO); + KokkosFFT::irfftn(execution_space(), x, out_o, KokkosFFT::Normalization::ORTHO); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_f, KokkosFFT::Normalization::FORWARD); + KokkosFFT::irfftn(execution_space(), x, out_f, KokkosFFT::Normalization::FORWARD); multiply(out_o, 1.0/sqrt(static_cast(n0 * n1 * n2))); multiply(out_f, 1.0/static_cast(n0 * n1 * n2)); @@ -1089,16 +1089,16 @@ void test_irfftn_3dfft_3dview() { using axes_type = KokkosFFT::axis_type<3>; Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out, axes_type{-3, -2, -1}); // default: KokkosFFT::Normalization::BACKWARD + KokkosFFT::irfftn(execution_space(), x, out, axes_type{-3, -2, -1}); // default: KokkosFFT::Normalization::BACKWARD Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_b, axes_type{-3, -2, -1}, KokkosFFT::Normalization::BACKWARD); + KokkosFFT::irfftn(execution_space(), x, out_b, axes_type{-3, -2, -1}, KokkosFFT::Normalization::BACKWARD); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_o, axes_type{-3, -2, -1}, KokkosFFT::Normalization::ORTHO); + KokkosFFT::irfftn(execution_space(), x, out_o, axes_type{-3, -2, -1}, KokkosFFT::Normalization::ORTHO); Kokkos::deep_copy(x, x_ref); - KokkosFFT::irfftn(x, out_f, axes_type{-3, -2, -1}, KokkosFFT::Normalization::FORWARD); + KokkosFFT::irfftn(execution_space(), x, out_f, axes_type{-3, -2, -1}, KokkosFFT::Normalization::FORWARD); multiply(out_o, 1.0/sqrt(static_cast(n0 * n1 * n2))); multiply(out_f, 1.0/static_cast(n0 * n1 * n2));