From 138ae122dba4c5052edb5b2730f954b7b748f54c Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Tue, 23 May 2023 22:13:23 -0400 Subject: [PATCH] Test CUDAMatrixFree not only for DEAL_II_WITH_CUDA --- .../deal.II/matrix_free/cuda_fe_evaluation.h | 24 +++--- .../matrix_free/cuda_hanging_nodes_internal.h | 10 +-- .../deal.II/matrix_free/cuda_matrix_free.h | 74 ++++++++----------- .../matrix_free/cuda_matrix_free.templates.h | 69 +++++++++-------- .../matrix_free/cuda_tensor_product_kernels.h | 4 - source/matrix_free/CMakeLists.txt | 11 +-- source/matrix_free/cuda_matrix_free.cc | 4 - .../matrix_free_device_initialize_vector.cc} | 8 +- ...with_mpi=on.with_p4est=on.mpirun=2.output} | 0 ...th_p4est=on.with_cuda=off.mpirun=1.output} | 5 -- ...with_p4est=on.with_cuda=on.mpirun=1.output | 13 ++++ .../matrix_free_device_matrix_vector_01.cc} | 4 +- ...atrix_free_device_matrix_vector_01.output} | 0 .../matrix_free_device_matrix_vector_02.cc} | 4 +- ...atrix_free_device_matrix_vector_02.output} | 0 .../matrix_free_device_matrix_vector_03.cc} | 4 +- ...atrix_free_device_matrix_vector_03.output} | 0 .../matrix_free_device_matrix_vector_03b.cc} | 4 +- ...trix_free_device_matrix_vector_03b.output} | 0 .../matrix_free_device_matrix_vector_04.cc} | 4 +- ...atrix_free_device_matrix_vector_04.output} | 0 .../matrix_free_device_matrix_vector_05.cc} | 4 +- ...atrix_free_device_matrix_vector_05.output} | 0 .../matrix_free_device_matrix_vector_06.cc} | 4 +- ...atrix_free_device_matrix_vector_06.output} | 0 .../matrix_free_device_matrix_vector_06a.cc} | 4 +- ...trix_free_device_matrix_vector_06a.output} | 0 .../matrix_free_device_matrix_vector_06b.cc} | 4 +- ...trix_free_device_matrix_vector_06b.output} | 0 .../matrix_free_device_matrix_vector_09.cc} | 4 +- ...atrix_free_device_matrix_vector_09.output} | 0 .../matrix_free_device_matrix_vector_10.cc} | 17 ++--- ...inos=true.with_p4est=true.mpirun=1.output} | 0 ...inos=true.with_p4est=true.mpirun=2.output} | 0 .../matrix_free_device_matrix_vector_10a.cc} | 17 ++--- ...inos=true.with_p4est=true.mpirun=1.output} | 0 ...inos=true.with_p4est=true.mpirun=2.output} | 0 .../matrix_free_device_matrix_vector_19.cc} | 17 ++--- ...trilinos_with_zoltan=true.mpirun=2.output} | 0 .../matrix_free_device_matrix_vector_25.cc} | 17 ++--- ...inos=true.with_p4est=true.mpirun=1.output} | 0 ...inos=true.with_p4est=true.mpirun=2.output} | 0 .../matrix_free_device_matrix_vector_26.cc} | 17 ++--- ...inos=true.with_p4est=true.mpirun=1.output} | 0 ...inos=true.with_p4est=true.mpirun=2.output} | 0 .../matrix_free_device_multiple_objects.cc} | 50 +++++++------ ...atrix_free_device_multiple_objects.output} | 0 ...matrix_free_device_no_index_initialize.cc} | 10 +-- ...ix_free_device_no_index_initialize.output} | 0 .../matrix_free_device_reinit_01.cc} | 1 - ...with_mpi=on.with_p4est=on.mpirun=2.output} | 0 .../matrix_vector_device_common.h} | 6 +- 52 files changed, 191 insertions(+), 223 deletions(-) rename tests/{cuda/matrix_free_initialize_vector.cc => matrix_free/matrix_free_device_initialize_vector.cc} (97%) rename tests/{cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=2.output => matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.mpirun=2.output} (100%) rename tests/{cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=1.output => matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=off.mpirun=1.output} (50%) create mode 100644 tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=on.mpirun=1.output rename tests/{cuda/matrix_free_matrix_vector_01.cc => matrix_free/matrix_free_device_matrix_vector_01.cc} (91%) rename tests/{cuda/matrix_free_matrix_vector_01.output => matrix_free/matrix_free_device_matrix_vector_01.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_02.cc => matrix_free/matrix_free_device_matrix_vector_02.cc} (93%) rename tests/{cuda/matrix_free_matrix_vector_02.output => matrix_free/matrix_free_device_matrix_vector_02.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_03.cc => matrix_free/matrix_free_device_matrix_vector_03.cc} (95%) rename tests/{cuda/matrix_free_matrix_vector_03.output => matrix_free/matrix_free_device_matrix_vector_03.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_03b.cc => matrix_free/matrix_free_device_matrix_vector_03b.cc} (95%) rename tests/{cuda/matrix_free_matrix_vector_03b.output => matrix_free/matrix_free_device_matrix_vector_03b.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_04.cc => matrix_free/matrix_free_device_matrix_vector_04.cc} (92%) rename tests/{cuda/matrix_free_matrix_vector_04.output => matrix_free/matrix_free_device_matrix_vector_04.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_05.cc => matrix_free/matrix_free_device_matrix_vector_05.cc} (95%) rename tests/{cuda/matrix_free_matrix_vector_05.output => matrix_free/matrix_free_device_matrix_vector_05.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_06.cc => matrix_free/matrix_free_device_matrix_vector_06.cc} (95%) rename tests/{cuda/matrix_free_matrix_vector_06.output => matrix_free/matrix_free_device_matrix_vector_06.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_06a.cc => matrix_free/matrix_free_device_matrix_vector_06a.cc} (97%) rename tests/{cuda/matrix_free_matrix_vector_06a.output => matrix_free/matrix_free_device_matrix_vector_06a.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_06b.cc => matrix_free/matrix_free_device_matrix_vector_06b.cc} (97%) rename tests/{cuda/matrix_free_matrix_vector_06b.output => matrix_free/matrix_free_device_matrix_vector_06b.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_09.cc => matrix_free/matrix_free_device_matrix_vector_09.cc} (95%) rename tests/{cuda/matrix_free_matrix_vector_09.output => matrix_free/matrix_free_device_matrix_vector_09.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_10.cc => matrix_free/matrix_free_device_matrix_vector_10.cc} (95%) rename tests/{cuda/matrix_free_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=1.output => matrix_free/matrix_free_device_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=1.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=2.output => matrix_free/matrix_free_device_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=2.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_10a.cc => matrix_free/matrix_free_device_matrix_vector_10a.cc} (95%) rename tests/{cuda/matrix_free_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=1.output => matrix_free/matrix_free_device_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=1.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=2.output => matrix_free/matrix_free_device_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=2.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_19.cc => matrix_free/matrix_free_device_matrix_vector_19.cc} (95%) rename tests/{cuda/matrix_free_matrix_vector_19.with_trilinos=true.with_trilinos_with_zoltan=true.mpirun=2.output => matrix_free/matrix_free_device_matrix_vector_19.with_trilinos=true.with_trilinos_with_zoltan=true.mpirun=2.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_25.cc => matrix_free/matrix_free_device_matrix_vector_25.cc} (95%) rename tests/{cuda/matrix_free_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=1.output => matrix_free/matrix_free_device_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=1.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=2.output => matrix_free/matrix_free_device_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=2.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_26.cc => matrix_free/matrix_free_device_matrix_vector_26.cc} (96%) rename tests/{cuda/matrix_free_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=1.output => matrix_free/matrix_free_device_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=1.output} (100%) rename tests/{cuda/matrix_free_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=2.output => matrix_free/matrix_free_device_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=2.output} (100%) rename tests/{cuda/matrix_free_multiple_objects.cc => matrix_free/matrix_free_device_multiple_objects.cc} (87%) rename tests/{cuda/matrix_free_multiple_objects.output => matrix_free/matrix_free_device_multiple_objects.output} (100%) rename tests/{cuda/matrix_free_no_index_initialize.cc => matrix_free/matrix_free_device_no_index_initialize.cc} (92%) rename tests/{cuda/matrix_free_no_index_initialize.output => matrix_free/matrix_free_device_no_index_initialize.output} (100%) rename tests/{cuda/matrix_free_reinit_01.cc => matrix_free/matrix_free_device_reinit_01.cc} (99%) rename tests/{cuda/matrix_free_reinit_01.with_mpi=on.with_p4est=on.mpirun=2.output => matrix_free/matrix_free_device_reinit_01.with_mpi=on.with_p4est=on.mpirun=2.output} (100%) rename tests/{cuda/matrix_vector_common.h => matrix_free/matrix_vector_device_common.h} (98%) diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h index b106f850b4..ae1b6ed9ff 100644 --- a/include/deal.II/matrix_free/cuda_fe_evaluation.h +++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h @@ -18,21 +18,19 @@ #include -#ifdef DEAL_II_WITH_CUDA +#include +#include +#include -# include -# include -# include +#include +#include -# include -# include +#include +#include +#include +#include -# include -# include -# include -# include - -# include +#include DEAL_II_NAMESPACE_OPEN @@ -526,5 +524,3 @@ namespace CUDAWrappers DEAL_II_NAMESPACE_CLOSE #endif - -#endif diff --git a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h index 90942f92fe..743b06f015 100644 --- a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h +++ b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h @@ -18,13 +18,11 @@ #include -#ifdef DEAL_II_WITH_CUDA +#include -# include +#include -# include - -# include +#include DEAL_II_NAMESPACE_OPEN namespace CUDAWrappers @@ -464,5 +462,3 @@ namespace CUDAWrappers DEAL_II_NAMESPACE_CLOSE #endif - -#endif diff --git a/include/deal.II/matrix_free/cuda_matrix_free.h b/include/deal.II/matrix_free/cuda_matrix_free.h index 3666d7e853..a30d57b260 100644 --- a/include/deal.II/matrix_free/cuda_matrix_free.h +++ b/include/deal.II/matrix_free/cuda_matrix_free.h @@ -19,28 +19,26 @@ #include -#ifdef DEAL_II_WITH_CUDA - -# include -# include -# include -# include -# include -# include -# include +#include +#include +#include +#include +#include +#include +#include -# include +#include -# include -# include +#include +#include -# include +#include -# include -# include -# include +#include +#include +#include -# include +#include @@ -58,13 +56,13 @@ namespace internal namespace CUDAWrappers { // forward declaration -# ifndef DOXYGEN +#ifndef DOXYGEN namespace internal { template class ReinitHelper; } -# endif +#endif /** * This class collects all the data that is stored for the matrix free @@ -118,12 +116,12 @@ namespace CUDAWrappers , use_coloring(use_coloring) , overlap_communication_computation(overlap_communication_computation) { -# ifndef DEAL_II_MPI_WITH_DEVICE_SUPPORT +#ifndef DEAL_II_MPI_WITH_DEVICE_SUPPORT AssertThrow( overlap_communication_computation == false, ExcMessage( "Overlapping communication and computation requires CUDA-aware MPI.")); -# endif +#endif if (overlap_communication_computation == true) AssertThrow( use_coloring == false || overlap_communication_computation == false, @@ -377,6 +375,7 @@ namespace CUDAWrappers void set_constrained_values(const Number val, VectorType &dst) const; +#ifdef DEAL_II_WITH_CUDA /** * Initialize a serial vector. The size corresponds to the number of degrees * of freedom in the DoFHandler object. @@ -384,6 +383,7 @@ namespace CUDAWrappers void initialize_dof_vector( LinearAlgebra::CUDAWrappers::Vector &vec) const; +#endif /** * Initialize a distributed vector. The local elements correspond to the @@ -392,7 +392,8 @@ namespace CUDAWrappers */ void initialize_dof_vector( - LinearAlgebra::distributed::Vector &vec) const; + LinearAlgebra::distributed::Vector &vec) + const; /** * Return the colored graph of locally owned active cells. @@ -456,10 +457,13 @@ namespace CUDAWrappers template void distributed_cell_loop( - const Functor & func, - const LinearAlgebra::distributed::Vector &src, - LinearAlgebra::distributed::Vector &dst) const; + const Functor &func, + const LinearAlgebra::distributed::Vector + & src, + LinearAlgebra::distributed::Vector &dst) + const; +#ifdef DEAL_II_WITH_CUDA /** * This function should never be called. Calling it results in an internal * error. This function exists only because cell_loop needs @@ -471,6 +475,7 @@ namespace CUDAWrappers const Functor & func, const LinearAlgebra::CUDAWrappers::Vector &src, LinearAlgebra::CUDAWrappers::Vector & dst) const; +#endif /** * Unique ID associated with the object. @@ -591,18 +596,6 @@ namespace CUDAWrappers Kokkos::View constraint_weights; - /** - * Grid dimensions associated to the different colors. The grid dimensions - * are used to launch the CUDA kernels. - */ - std::vector grid_dim; - - /** - * Block dimensions associated to the different colors. The block dimensions - * are used to launch the CUDA kernels. - */ - std::vector block_dim; - /** * Shared pointer to a Partitioner for distributed Vectors used in * cell_loop. When MPI is not used the pointer is null. @@ -788,8 +781,6 @@ namespace CUDAWrappers data_host.row_start = data.row_start; data_host.use_coloring = data.use_coloring; - const unsigned int n_elements = - data_host.n_cells * data_host.padding_length; if (update_flags & update_quadrature_points) { data_host.q_points = Kokkos::create_mirror(data.q_points); @@ -820,7 +811,7 @@ namespace CUDAWrappers /*----------------------- Inline functions ---------------------------------*/ -# ifndef DOXYGEN +#ifndef DOXYGEN template inline const std::vector -#ifdef DEAL_II_WITH_CUDA - -# include -# include -# include -# include +#include +#include +#include +#include -# include +#include -# include -# include -# include +#include +#include +#include -# include -# include -# include +#include +#include +#include -# include +#include +#ifdef DEAL_II_WITH_CUDA # include +#endif -# include -# include -# include +#include +#include +#include DEAL_II_NAMESPACE_OPEN @@ -132,8 +132,6 @@ namespace CUDAWrappers // We need at least three colors when we are using CUDA-aware MPI and // overlapping the communication data->n_cells.resize(std::max(n_colors, 3U), 0); - data->grid_dim.resize(n_colors); - data->block_dim.resize(n_colors); data->local_to_global.resize(n_colors); data->constraint_mask.resize(n_colors); @@ -309,6 +307,7 @@ namespace CUDAWrappers } }; +#ifdef DEAL_II_WITH_CUDA template <> struct VectorLocalSize> { @@ -328,6 +327,7 @@ namespace CUDAWrappers return vec.size(); } }; +#endif @@ -557,7 +557,7 @@ namespace CUDAWrappers } - +#ifdef DEAL_II_WITH_CUDA template void MatrixFree::initialize_dof_vector( @@ -565,13 +565,14 @@ namespace CUDAWrappers { vec.reinit(n_dofs); } +#endif template void MatrixFree::initialize_dof_vector( - LinearAlgebra::distributed::Vector &vec) const + LinearAlgebra::distributed::Vector &vec) const { if (partitioner) vec.reinit(partitioner); @@ -635,7 +636,6 @@ namespace CUDAWrappers // First compute the size of n_cells, row_starts, kernel launch parameters, // and constrained_dofs std::size_t bytes = n_cells.size() * sizeof(unsigned int) * 2 + - 2 * n_colors * sizeof(dim3) + n_constrained_dofs * sizeof(unsigned int); // For each color, add local_to_global, inv_jacobian, JxW, and q_points. @@ -668,9 +668,6 @@ namespace CUDAWrappers { dof_handler = &dof_handler_; - if (typeid(Number) == typeid(double)) - cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte); - UpdateFlags update_flags = additional_data.mapping_update_flags; if (update_flags & update_gradients) update_flags |= update_JxW_values; @@ -937,9 +934,9 @@ namespace CUDAWrappers template void MatrixFree::distributed_cell_loop( - const Functor & func, - const LinearAlgebra::distributed::Vector &src, - LinearAlgebra::distributed::Vector &dst) const + const Functor &func, + const LinearAlgebra::distributed::Vector &src, + LinearAlgebra::distributed::Vector &dst) const { MemorySpace::Default::kokkos_space::execution_space exec; @@ -1046,12 +1043,13 @@ namespace CUDAWrappers else { // Create the ghosted source and the ghosted destination - LinearAlgebra::distributed::Vector + LinearAlgebra::distributed::Vector ghosted_src(partitioner); - LinearAlgebra::distributed::Vector + LinearAlgebra::distributed::Vector ghosted_dst(ghosted_src); ghosted_src = src; ghosted_dst = dst; + ghosted_dst.zero_out_ghost_values(); // Execute the loop on the cells for (unsigned int i = 0; i < n_colors; ++i) @@ -1062,7 +1060,10 @@ namespace CUDAWrappers team_policy(exec, n_cells[i], Kokkos::AUTO); internal::ApplyKernel apply_kernel( - func, get_data(i), src.get_values(), dst.get_values()); + func, + get_data(i), + ghosted_src.get_values(), + ghosted_dst.get_values()); Kokkos::parallel_for( "dealii::MatrixFree::distributed_cell_loop_" + @@ -1077,8 +1078,7 @@ namespace CUDAWrappers } } - - +#ifdef DEAL_II_WITH_CUDA template template void @@ -1089,10 +1089,9 @@ namespace CUDAWrappers { Assert(false, ExcInternalError()); } +#endif } // namespace CUDAWrappers DEAL_II_NAMESPACE_CLOSE #endif - -#endif diff --git a/include/deal.II/matrix_free/cuda_tensor_product_kernels.h b/include/deal.II/matrix_free/cuda_tensor_product_kernels.h index b191a0cf5b..ae39a56f51 100644 --- a/include/deal.II/matrix_free/cuda_tensor_product_kernels.h +++ b/include/deal.II/matrix_free/cuda_tensor_product_kernels.h @@ -23,8 +23,6 @@ #include -#ifdef DEAL_II_WITH_CUDA - DEAL_II_NAMESPACE_OPEN @@ -756,5 +754,3 @@ namespace CUDAWrappers DEAL_II_NAMESPACE_CLOSE #endif - -#endif diff --git a/source/matrix_free/CMakeLists.txt b/source/matrix_free/CMakeLists.txt index 995f85a391..19212afd94 100644 --- a/source/matrix_free/CMakeLists.txt +++ b/source/matrix_free/CMakeLists.txt @@ -14,6 +14,7 @@ ## --------------------------------------------------------------------- set(_src + cuda_matrix_free.cc dof_info.cc evaluation_template_factory.cc evaluation_template_factory_inst2.cc @@ -52,15 +53,5 @@ file(GLOB _header ${CMAKE_SOURCE_DIR}/include/deal.II/matrix_free/*.h ) -if(DEAL_II_WITH_CUDA) - set(_src - cuda_matrix_free.cc - ${_src} - ) - set(_inst - ${_inst} - ) -endif() - define_object_library(object_matrix_free OBJECT ${_src} ${_header} ${_inst}) expand_instantiations(object_matrix_free "${_inst}") diff --git a/source/matrix_free/cuda_matrix_free.cc b/source/matrix_free/cuda_matrix_free.cc index 2ff0a5968f..22f89f3b11 100644 --- a/source/matrix_free/cuda_matrix_free.cc +++ b/source/matrix_free/cuda_matrix_free.cc @@ -15,8 +15,6 @@ #include -#ifdef DEAL_II_WITH_CUDA - DEAL_II_NAMESPACE_OPEN @@ -31,5 +29,3 @@ namespace CUDAWrappers } // namespace CUDAWrappers DEAL_II_NAMESPACE_CLOSE - -#endif diff --git a/tests/cuda/matrix_free_initialize_vector.cc b/tests/matrix_free/matrix_free_device_initialize_vector.cc similarity index 97% rename from tests/cuda/matrix_free_initialize_vector.cc rename to tests/matrix_free/matrix_free_device_initialize_vector.cc index 7cdefebe5f..0c53f24828 100644 --- a/tests/cuda/matrix_free_initialize_vector.cc +++ b/tests/matrix_free/matrix_free_device_initialize_vector.cc @@ -50,6 +50,7 @@ check(const LinearAlgebra::distributed::Vector ExcInternalError()); } +#ifdef DEAL_II_WITH_CUDA template void check(const LinearAlgebra::CUDAWrappers::Vector &vector, @@ -57,6 +58,7 @@ check(const LinearAlgebra::CUDAWrappers::Vector &vector, { AssertDimension(vector.size(), reference_partitioner.size()); } +#endif template void @@ -113,14 +115,16 @@ main(int argc, char **argv) unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD); deallog.push(Utilities::int_to_string(myid)); - init_cuda(true); + // init_cuda(true); MPILogInitAll mpi_inilog; test<2, 1, LinearAlgebra::distributed::Vector>(); + +#ifdef DEAL_II_WITH_CUDA if (Utilities::MPI::n_mpi_processes(MPI_COMM_WORLD) == 1) test<2, 1, LinearAlgebra::CUDAWrappers::Vector>(); - +#endif return 0; } diff --git a/tests/cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=2.output b/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.mpirun=2.output similarity index 100% rename from tests/cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=2.output rename to tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.mpirun=2.output diff --git a/tests/cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=1.output b/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=off.mpirun=1.output similarity index 50% rename from tests/cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=1.output rename to tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=off.mpirun=1.output index a0d4178229..274361dac4 100644 --- a/tests/cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=1.output +++ b/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=off.mpirun=1.output @@ -4,8 +4,3 @@ DEAL:0:0::locally owned dofs : DEAL:0:0::locally relevant dofs : {[0,24]} DEAL:0:0::OK -DEAL:0:0::locally owned dofs : -{[0,24]} -DEAL:0:0::locally relevant dofs : -{[0,24]} -DEAL:0:0::OK diff --git a/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=on.mpirun=1.output b/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=on.mpirun=1.output new file mode 100644 index 0000000000..7e4be316a7 --- /dev/null +++ b/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=on.mpirun=1.output @@ -0,0 +1,13 @@ + + +- +-DEAL:0:0::locally owned dofs : +-{[0,24]} +-DEAL:0:0::locally relevant dofs : +-{[0,24]} +-DEAL:0:0::OK +-DEAL:0:0::locally owned dofs : +-{[0,24]} +-DEAL:0:0::locally relevant dofs : +-{[0,24]} +-DEAL:0:0::OK diff --git a/tests/cuda/matrix_free_matrix_vector_01.cc b/tests/matrix_free/matrix_free_device_matrix_vector_01.cc similarity index 91% rename from tests/cuda/matrix_free_matrix_vector_01.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_01.cc index c6a79083fc..e626ae812c 100644 --- a/tests/cuda/matrix_free_matrix_vector_01.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_01.cc @@ -22,7 +22,7 @@ #include "../tests.h" -#include "matrix_vector_common.h" +#include "matrix_vector_device_common.h" template void @@ -41,6 +41,6 @@ test() do_test, + LinearAlgebra::distributed::Vector, fe_degree + 1>(dof, constraints, tria.n_active_cells()); } diff --git a/tests/cuda/matrix_free_matrix_vector_01.output b/tests/matrix_free/matrix_free_device_matrix_vector_01.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_01.output rename to tests/matrix_free/matrix_free_device_matrix_vector_01.output diff --git a/tests/cuda/matrix_free_matrix_vector_02.cc b/tests/matrix_free/matrix_free_device_matrix_vector_02.cc similarity index 93% rename from tests/cuda/matrix_free_matrix_vector_02.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_02.cc index ceaa73a31a..f50cb6c522 100644 --- a/tests/cuda/matrix_free_matrix_vector_02.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_02.cc @@ -22,7 +22,7 @@ #include "../tests.h" -#include "matrix_vector_common.h" +#include "matrix_vector_device_common.h" template void @@ -45,6 +45,6 @@ test() do_test, + LinearAlgebra::distributed::Vector, fe_degree + 1>(dof, constraints, tria.n_active_cells()); } diff --git a/tests/cuda/matrix_free_matrix_vector_02.output b/tests/matrix_free/matrix_free_device_matrix_vector_02.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_02.output rename to tests/matrix_free/matrix_free_device_matrix_vector_02.output diff --git a/tests/cuda/matrix_free_matrix_vector_03.cc b/tests/matrix_free/matrix_free_device_matrix_vector_03.cc similarity index 95% rename from tests/cuda/matrix_free_matrix_vector_03.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_03.cc index cd581a99c5..2894340c29 100644 --- a/tests/cuda/matrix_free_matrix_vector_03.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_03.cc @@ -24,7 +24,7 @@ #include "../tests.h" -#include "matrix_vector_common.h" +#include "matrix_vector_device_common.h" template void @@ -66,6 +66,6 @@ test() do_test, + LinearAlgebra::distributed::Vector, fe_degree + 1>(dof, constraints, tria.n_active_cells()); } diff --git a/tests/cuda/matrix_free_matrix_vector_03.output b/tests/matrix_free/matrix_free_device_matrix_vector_03.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_03.output rename to tests/matrix_free/matrix_free_device_matrix_vector_03.output diff --git a/tests/cuda/matrix_free_matrix_vector_03b.cc b/tests/matrix_free/matrix_free_device_matrix_vector_03b.cc similarity index 95% rename from tests/cuda/matrix_free_matrix_vector_03b.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_03b.cc index 464aadc676..38e924ea14 100644 --- a/tests/cuda/matrix_free_matrix_vector_03b.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_03b.cc @@ -25,7 +25,7 @@ #include "../tests.h" -#include "matrix_vector_common.h" +#include "matrix_vector_device_common.h" template void @@ -68,6 +68,6 @@ test() do_test, + LinearAlgebra::distributed::Vector, fe_degree + 1>(dof, constraints, tria.n_active_cells(), true, true); } diff --git a/tests/cuda/matrix_free_matrix_vector_03b.output b/tests/matrix_free/matrix_free_device_matrix_vector_03b.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_03b.output rename to tests/matrix_free/matrix_free_device_matrix_vector_03b.output diff --git a/tests/cuda/matrix_free_matrix_vector_04.cc b/tests/matrix_free/matrix_free_device_matrix_vector_04.cc similarity index 92% rename from tests/cuda/matrix_free_matrix_vector_04.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_04.cc index c27adc0907..f68d81e5fb 100644 --- a/tests/cuda/matrix_free_matrix_vector_04.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_04.cc @@ -21,7 +21,7 @@ #include "../tests.h" -#include "matrix_vector_common.h" +#include "matrix_vector_device_common.h" template @@ -45,6 +45,6 @@ test() do_test, + LinearAlgebra::distributed::Vector, fe_degree + 1>(dof, constraints, tria.n_active_cells()); } diff --git a/tests/cuda/matrix_free_matrix_vector_04.output b/tests/matrix_free/matrix_free_device_matrix_vector_04.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_04.output rename to tests/matrix_free/matrix_free_device_matrix_vector_04.output diff --git a/tests/cuda/matrix_free_matrix_vector_05.cc b/tests/matrix_free/matrix_free_device_matrix_vector_05.cc similarity index 95% rename from tests/cuda/matrix_free_matrix_vector_05.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_05.cc index cffa01115f..b8daf713cc 100644 --- a/tests/cuda/matrix_free_matrix_vector_05.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_05.cc @@ -22,7 +22,7 @@ #include "../tests.h" -#include "matrix_vector_common.h" +#include "matrix_vector_device_common.h" template @@ -77,6 +77,6 @@ test() do_test, + LinearAlgebra::distributed::Vector, fe_degree + 1>(dof, constraints, tria.n_active_cells()); } diff --git a/tests/cuda/matrix_free_matrix_vector_05.output b/tests/matrix_free/matrix_free_device_matrix_vector_05.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_05.output rename to tests/matrix_free/matrix_free_device_matrix_vector_05.output diff --git a/tests/cuda/matrix_free_matrix_vector_06.cc b/tests/matrix_free/matrix_free_device_matrix_vector_06.cc similarity index 95% rename from tests/cuda/matrix_free_matrix_vector_06.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_06.cc index 27fdf46f73..14776e9dc8 100644 --- a/tests/cuda/matrix_free_matrix_vector_06.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_06.cc @@ -26,7 +26,7 @@ #include "../tests.h" #include "create_mesh.h" -#include "matrix_vector_common.h" +#include "matrix_vector_device_common.h" template void @@ -75,6 +75,6 @@ test() do_test, + LinearAlgebra::distributed::Vector, fe_degree + 1>(dof, constraints, tria.n_active_cells()); } diff --git a/tests/cuda/matrix_free_matrix_vector_06.output b/tests/matrix_free/matrix_free_device_matrix_vector_06.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_06.output rename to tests/matrix_free/matrix_free_device_matrix_vector_06.output diff --git a/tests/cuda/matrix_free_matrix_vector_06a.cc b/tests/matrix_free/matrix_free_device_matrix_vector_06a.cc similarity index 97% rename from tests/cuda/matrix_free_matrix_vector_06a.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_06a.cc index e20b65caaa..41ea3e611b 100644 --- a/tests/cuda/matrix_free_matrix_vector_06a.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_06a.cc @@ -22,7 +22,7 @@ #include "../tests.h" #include "create_mesh.h" -#include "matrix_vector_common.h" +#include "matrix_vector_device_common.h" template void @@ -71,6 +71,6 @@ test() do_test, + LinearAlgebra::distributed::Vector, fe_degree + 1>(dof, constraints, tria.n_active_cells()); } diff --git a/tests/cuda/matrix_free_matrix_vector_06a.output b/tests/matrix_free/matrix_free_device_matrix_vector_06a.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_06a.output rename to tests/matrix_free/matrix_free_device_matrix_vector_06a.output diff --git a/tests/cuda/matrix_free_matrix_vector_06b.cc b/tests/matrix_free/matrix_free_device_matrix_vector_06b.cc similarity index 97% rename from tests/cuda/matrix_free_matrix_vector_06b.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_06b.cc index 473297c9b3..5d7d72718a 100644 --- a/tests/cuda/matrix_free_matrix_vector_06b.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_06b.cc @@ -21,7 +21,7 @@ #include "../tests.h" #include "create_mesh.h" -#include "matrix_vector_common.h" +#include "matrix_vector_device_common.h" template void @@ -70,6 +70,6 @@ test() do_test, + LinearAlgebra::distributed::Vector, fe_degree + 1>(dof, constraints, tria.n_active_cells(), false); } diff --git a/tests/cuda/matrix_free_matrix_vector_06b.output b/tests/matrix_free/matrix_free_device_matrix_vector_06b.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_06b.output rename to tests/matrix_free/matrix_free_device_matrix_vector_06b.output diff --git a/tests/cuda/matrix_free_matrix_vector_09.cc b/tests/matrix_free/matrix_free_device_matrix_vector_09.cc similarity index 95% rename from tests/cuda/matrix_free_matrix_vector_09.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_09.cc index 13244c274d..833319faef 100644 --- a/tests/cuda/matrix_free_matrix_vector_09.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_09.cc @@ -24,7 +24,7 @@ #include "../tests.h" #include "create_mesh.h" -#include "matrix_vector_common.h" +#include "matrix_vector_device_common.h" template void @@ -73,6 +73,6 @@ test() do_test, + LinearAlgebra::distributed::Vector, fe_degree + 1>(dof, constraints, tria.n_active_cells()); } diff --git a/tests/cuda/matrix_free_matrix_vector_09.output b/tests/matrix_free/matrix_free_device_matrix_vector_09.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_09.output rename to tests/matrix_free/matrix_free_device_matrix_vector_09.output diff --git a/tests/cuda/matrix_free_matrix_vector_10.cc b/tests/matrix_free/matrix_free_device_matrix_vector_10.cc similarity index 95% rename from tests/cuda/matrix_free_matrix_vector_10.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_10.cc index 6ea858759a..92b50ded61 100644 --- a/tests/cuda/matrix_free_matrix_vector_10.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_10.cc @@ -46,7 +46,7 @@ #include "../tests.h" -#include "matrix_vector_mf.h" +#include "matrix_vector_device_mf.h" @@ -117,13 +117,14 @@ test() const unsigned int coef_size = tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim); - MatrixFreeTest> + MatrixFreeTest< + dim, + fe_degree, + Number, + LinearAlgebra::distributed::Vector> mf(mf_data, coef_size); - LinearAlgebra::distributed::Vector in_dev; - LinearAlgebra::distributed::Vector out_dev; + LinearAlgebra::distributed::Vector in_dev; + LinearAlgebra::distributed::Vector out_dev; mf_data.initialize_dof_vector(in_dev); mf_data.initialize_dof_vector(out_dev); @@ -224,8 +225,6 @@ main(int argc, char **argv) unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD); deallog.push(Utilities::int_to_string(myid)); - init_cuda(true); - if (myid == 0) { initlog(); diff --git a/tests/cuda/matrix_free_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=1.output b/tests/matrix_free/matrix_free_device_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=1.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=1.output rename to tests/matrix_free/matrix_free_device_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=1.output diff --git a/tests/cuda/matrix_free_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=2.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=2.output rename to tests/matrix_free/matrix_free_device_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=2.output diff --git a/tests/cuda/matrix_free_matrix_vector_10a.cc b/tests/matrix_free/matrix_free_device_matrix_vector_10a.cc similarity index 95% rename from tests/cuda/matrix_free_matrix_vector_10a.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_10a.cc index 7c434dd7bb..513c8bdb82 100644 --- a/tests/cuda/matrix_free_matrix_vector_10a.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_10a.cc @@ -41,7 +41,7 @@ #include "../tests.h" -#include "matrix_vector_mf.h" +#include "matrix_vector_device_mf.h" template @@ -119,14 +119,15 @@ test() const unsigned int coef_size = tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim); - MatrixFreeTest> + MatrixFreeTest< + dim, + fe_degree, + Number, + LinearAlgebra::distributed::Vector> mf(mf_data, coef_size, false); - LinearAlgebra::distributed::Vector in_dev; - LinearAlgebra::distributed::Vector out_dev; + LinearAlgebra::distributed::Vector in_dev; + LinearAlgebra::distributed::Vector out_dev; mf_data.initialize_dof_vector(in_dev); mf_data.initialize_dof_vector(out_dev); LinearAlgebra::ReadWriteVector rw_in(owned_set); @@ -230,8 +231,6 @@ main(int argc, char **argv) unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD); deallog.push(Utilities::int_to_string(myid)); - init_cuda(true); - if (myid == 0) { initlog(); diff --git a/tests/cuda/matrix_free_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=1.output b/tests/matrix_free/matrix_free_device_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=1.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=1.output rename to tests/matrix_free/matrix_free_device_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=1.output diff --git a/tests/cuda/matrix_free_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=2.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=2.output rename to tests/matrix_free/matrix_free_device_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=2.output diff --git a/tests/cuda/matrix_free_matrix_vector_19.cc b/tests/matrix_free/matrix_free_device_matrix_vector_19.cc similarity index 95% rename from tests/cuda/matrix_free_matrix_vector_19.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_19.cc index 39fba64516..f71c3c5f33 100644 --- a/tests/cuda/matrix_free_matrix_vector_19.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_19.cc @@ -42,7 +42,7 @@ #include "../tests.h" -#include "matrix_vector_mf.h" +#include "matrix_vector_device_mf.h" @@ -113,14 +113,15 @@ test() const unsigned int coef_size = tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim); - MatrixFreeTest> + MatrixFreeTest< + dim, + fe_degree, + Number, + LinearAlgebra::distributed::Vector> mf(mf_data, coef_size); - LinearAlgebra::distributed::Vector in_dev( + LinearAlgebra::distributed::Vector in_dev( owned_set, MPI_COMM_WORLD); - LinearAlgebra::distributed::Vector out_dev( + LinearAlgebra::distributed::Vector out_dev( owned_set, MPI_COMM_WORLD); LinearAlgebra::ReadWriteVector rw_in(owned_set); @@ -220,8 +221,6 @@ main(int argc, char **argv) unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD); deallog.push(Utilities::int_to_string(myid)); - init_cuda(true); - if (myid == 0) { initlog(); diff --git a/tests/cuda/matrix_free_matrix_vector_19.with_trilinos=true.with_trilinos_with_zoltan=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_matrix_vector_19.with_trilinos=true.with_trilinos_with_zoltan=true.mpirun=2.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_19.with_trilinos=true.with_trilinos_with_zoltan=true.mpirun=2.output rename to tests/matrix_free/matrix_free_device_matrix_vector_19.with_trilinos=true.with_trilinos_with_zoltan=true.mpirun=2.output diff --git a/tests/cuda/matrix_free_matrix_vector_25.cc b/tests/matrix_free/matrix_free_device_matrix_vector_25.cc similarity index 95% rename from tests/cuda/matrix_free_matrix_vector_25.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_25.cc index a80c6ac3b4..dcb5464c43 100644 --- a/tests/cuda/matrix_free_matrix_vector_25.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_25.cc @@ -40,7 +40,7 @@ #include "../tests.h" -#include "matrix_vector_mf.h" +#include "matrix_vector_device_mf.h" @@ -120,13 +120,14 @@ test() const unsigned int coef_size = tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim); - MatrixFreeTest> + MatrixFreeTest< + dim, + fe_degree, + Number, + LinearAlgebra::distributed::Vector> mf(mf_data, coef_size); - LinearAlgebra::distributed::Vector in_dev; - LinearAlgebra::distributed::Vector out_dev; + LinearAlgebra::distributed::Vector in_dev; + LinearAlgebra::distributed::Vector out_dev; mf_data.initialize_dof_vector(in_dev); mf_data.initialize_dof_vector(out_dev); @@ -227,8 +228,6 @@ main(int argc, char **argv) unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD); deallog.push(Utilities::int_to_string(myid)); - init_cuda(true); - if (myid == 0) { initlog(); diff --git a/tests/cuda/matrix_free_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=1.output b/tests/matrix_free/matrix_free_device_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=1.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=1.output rename to tests/matrix_free/matrix_free_device_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=1.output diff --git a/tests/cuda/matrix_free_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=2.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=2.output rename to tests/matrix_free/matrix_free_device_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=2.output diff --git a/tests/cuda/matrix_free_matrix_vector_26.cc b/tests/matrix_free/matrix_free_device_matrix_vector_26.cc similarity index 96% rename from tests/cuda/matrix_free_matrix_vector_26.cc rename to tests/matrix_free/matrix_free_device_matrix_vector_26.cc index 8c1eeaf803..32b73b449b 100644 --- a/tests/cuda/matrix_free_matrix_vector_26.cc +++ b/tests/matrix_free/matrix_free_device_matrix_vector_26.cc @@ -42,7 +42,7 @@ #include "../tests.h" -#include "matrix_vector_mf.h" +#include "matrix_vector_device_mf.h" @@ -126,13 +126,14 @@ test() const unsigned int coef_size = tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim); - MatrixFreeTest> + MatrixFreeTest< + dim, + fe_degree, + Number, + LinearAlgebra::distributed::Vector> mf(mf_data, coef_size); - LinearAlgebra::distributed::Vector in_dev; - LinearAlgebra::distributed::Vector out_dev; + LinearAlgebra::distributed::Vector in_dev; + LinearAlgebra::distributed::Vector out_dev; mf_data.initialize_dof_vector(in_dev); mf_data.initialize_dof_vector(out_dev); @@ -233,8 +234,6 @@ main(int argc, char **argv) unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD); deallog.push(Utilities::int_to_string(myid)); - init_cuda(true); - if (myid == 0) { initlog(); diff --git a/tests/cuda/matrix_free_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=1.output b/tests/matrix_free/matrix_free_device_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=1.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=1.output rename to tests/matrix_free/matrix_free_device_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=1.output diff --git a/tests/cuda/matrix_free_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=2.output similarity index 100% rename from tests/cuda/matrix_free_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=2.output rename to tests/matrix_free/matrix_free_device_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=2.output diff --git a/tests/cuda/matrix_free_multiple_objects.cc b/tests/matrix_free/matrix_free_device_multiple_objects.cc similarity index 87% rename from tests/cuda/matrix_free_multiple_objects.cc rename to tests/matrix_free/matrix_free_device_multiple_objects.cc index e1bf9da3b2..04d44ba029 100644 --- a/tests/cuda/matrix_free_multiple_objects.cc +++ b/tests/matrix_free/matrix_free_device_multiple_objects.cc @@ -44,24 +44,27 @@ #include "../tests.h" #include "Kokkos_Core.hpp" -#include "matrix_vector_mf.h" +#include "matrix_vector_device_mf.h" template void -do_test(const DoFHandler<2> & dof, - MatrixFreeTest<2, - fe_degree, - double, - LinearAlgebra::CUDAWrappers::Vector, - n_q_points_1d> & mf, +do_test(const DoFHandler<2> &dof, + MatrixFreeTest< + 2, + fe_degree, + double, + LinearAlgebra::distributed::Vector, + n_q_points_1d> & mf, unsigned int n_dofs, MappingQ<2> & mapping, const AffineConstraints &constraints) { - Vector in_host(n_dofs), out_host(n_dofs); - LinearAlgebra::ReadWriteVector in(n_dofs), out(n_dofs); - LinearAlgebra::CUDAWrappers::Vector in_device(n_dofs); - LinearAlgebra::CUDAWrappers::Vector out_device(n_dofs); + Vector in_host(n_dofs), out_host(n_dofs); + LinearAlgebra::ReadWriteVector in(n_dofs), out(n_dofs); + LinearAlgebra::distributed::Vector in_device( + n_dofs); + LinearAlgebra::distributed::Vector out_device( + n_dofs); for (unsigned int i = 0; i < n_dofs; ++i) { @@ -73,7 +76,7 @@ do_test(const DoFHandler<2> & dof, in_device.import_elements(in, VectorOperation::insert); mf.vmult(out_device, in_device); - cudaDeviceSynchronize(); + Kokkos::fence(); out.import_elements(out_device, VectorOperation::insert); // assemble sparse matrix with (\nabla v, \nabla u) + (v, 10 * u) @@ -149,7 +152,6 @@ main() deallog << std::setprecision(3); Kokkos::initialize(); - init_cuda(); { Triangulation<2> tria; @@ -174,11 +176,12 @@ main() const QGauss<1> quad_1(n_q_points_1d_1); mf_data_1.reinit(mapping_1, dof_1, constraints, quad_1, additional_data_1); const unsigned int n_dofs_1 = dof_1.n_dofs(); - MatrixFreeTest<2, - fe_degree_1, - double, - LinearAlgebra::CUDAWrappers::Vector, - n_q_points_1d_1> + MatrixFreeTest< + 2, + fe_degree_1, + double, + LinearAlgebra::distributed::Vector, + n_q_points_1d_1> mf_1(mf_data_1, n_dofs_1 * std::pow(n_q_points_1d_1, 2), constant_coefficient); @@ -198,11 +201,12 @@ main() const QGauss<1> quad_2(n_q_points_1d_2); mf_data_2.reinit(mapping_2, dof_2, constraints, quad_2, additional_data_2); const unsigned int n_dofs_2 = dof_2.n_dofs(); - MatrixFreeTest<2, - fe_degree_2, - double, - LinearAlgebra::CUDAWrappers::Vector, - n_q_points_1d_2> + MatrixFreeTest< + 2, + fe_degree_2, + double, + LinearAlgebra::distributed::Vector, + n_q_points_1d_2> mf_2(mf_data_2, n_dofs_2 * std::pow(n_q_points_1d_2, 2), constant_coefficient); diff --git a/tests/cuda/matrix_free_multiple_objects.output b/tests/matrix_free/matrix_free_device_multiple_objects.output similarity index 100% rename from tests/cuda/matrix_free_multiple_objects.output rename to tests/matrix_free/matrix_free_device_multiple_objects.output diff --git a/tests/cuda/matrix_free_no_index_initialize.cc b/tests/matrix_free/matrix_free_device_no_index_initialize.cc similarity index 92% rename from tests/cuda/matrix_free_no_index_initialize.cc rename to tests/matrix_free/matrix_free_device_no_index_initialize.cc index bda055ddc6..9ab2222c4f 100644 --- a/tests/cuda/matrix_free_no_index_initialize.cc +++ b/tests/matrix_free/matrix_free_device_no_index_initialize.cc @@ -83,15 +83,12 @@ public: void test() const { - LinearAlgebra::distributed::Vector dst_dummy; - LinearAlgebra::distributed::Vector src_dummy; + LinearAlgebra::distributed::Vector dst_dummy; + LinearAlgebra::distributed::Vector src_dummy; data.cell_loop(*this, src_dummy, dst_dummy); - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); + Kokkos::fence(); deallog << "OK" << std::endl; }; @@ -160,7 +157,6 @@ main() initlog(); Kokkos::initialize(); - init_cuda(); test<2, 1>(); diff --git a/tests/cuda/matrix_free_no_index_initialize.output b/tests/matrix_free/matrix_free_device_no_index_initialize.output similarity index 100% rename from tests/cuda/matrix_free_no_index_initialize.output rename to tests/matrix_free/matrix_free_device_no_index_initialize.output diff --git a/tests/cuda/matrix_free_reinit_01.cc b/tests/matrix_free/matrix_free_device_reinit_01.cc similarity index 99% rename from tests/cuda/matrix_free_reinit_01.cc rename to tests/matrix_free/matrix_free_device_reinit_01.cc index 0d072baba3..6067986f16 100644 --- a/tests/cuda/matrix_free_reinit_01.cc +++ b/tests/matrix_free/matrix_free_device_reinit_01.cc @@ -91,7 +91,6 @@ main(int argc, char **argv) unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD); deallog.push(Utilities::int_to_string(myid)); - init_cuda(true); MPILogInitAll mpi_inilog; test<2, 1>(); diff --git a/tests/cuda/matrix_free_reinit_01.with_mpi=on.with_p4est=on.mpirun=2.output b/tests/matrix_free/matrix_free_device_reinit_01.with_mpi=on.with_p4est=on.mpirun=2.output similarity index 100% rename from tests/cuda/matrix_free_reinit_01.with_mpi=on.with_p4est=on.mpirun=2.output rename to tests/matrix_free/matrix_free_device_reinit_01.with_mpi=on.with_p4est=on.mpirun=2.output diff --git a/tests/cuda/matrix_vector_common.h b/tests/matrix_free/matrix_vector_device_common.h similarity index 98% rename from tests/cuda/matrix_vector_common.h rename to tests/matrix_free/matrix_vector_device_common.h index 7420cf5032..de11174444 100644 --- a/tests/cuda/matrix_vector_common.h +++ b/tests/matrix_free/matrix_vector_device_common.h @@ -47,7 +47,7 @@ #include "../tests.h" #include "Kokkos_Core.hpp" -#include "matrix_vector_mf.h" +#include "matrix_vector_device_mf.h" // forward declare this function. will be implemented in .cc files @@ -116,7 +116,7 @@ do_test(const DoFHandler & dof, in_device.import_elements(in, VectorOperation::insert); mf.vmult(out_device, in_device); - cudaDeviceSynchronize(); + Kokkos::fence(); out.import_elements(out_device, VectorOperation::insert); // assemble sparse matrix with (\nabla v, \nabla u) + (v, 10 * u) @@ -197,8 +197,6 @@ main() Kokkos::initialize(); - init_cuda(); - { deallog.push("double"); deallog.push("2d"); -- 2.39.5