From 21fed9d50f3e7a053156bd600c445e5284190a96 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Tue, 6 Dec 2022 11:23:54 -0500 Subject: [PATCH] Fix up --- include/deal.II/base/array_view.h | 48 +- .../lac/la_parallel_vector.templates.h | 108 ++-- .../deal.II/lac/vector_operations_internal.h | 503 +++++++++--------- source/lac/CMakeLists.txt | 2 +- source/lac/vector_memory.cc | 12 +- 5 files changed, 325 insertions(+), 348 deletions(-) diff --git a/include/deal.II/base/array_view.h b/include/deal.II/base/array_view.h index 50053b5330..168310404e 100644 --- a/include/deal.II/base/array_view.h +++ b/include/deal.II/base/array_view.h @@ -345,45 +345,6 @@ private: //--------------------------------------------------------------------------- -namespace internal -{ - namespace ArrayViewHelper - { - template - inline bool - is_in_correct_memory_space(const void *const ptr) - { -#ifndef DEAL_II_COMPILER_CUDA_AWARE - (void)ptr; - static_assert(std::is_same::value, - "If the compiler doesn't understand CUDA code, " - "the only possible memory space is 'MemorySpace::Host'!"); - return true; -#else - cudaPointerAttributes attributes; - const cudaError_t cuda_error = cudaPointerGetAttributes(&attributes, ptr); - if (cuda_error != cudaErrorInvalidValue) - { - AssertCuda(cuda_error); - if (std::is_same::value) - return (attributes.type == cudaMemoryTypeHost) || - (attributes.type == cudaMemoryTypeUnregistered); - else - return attributes.type == cudaMemoryTypeDevice; - } - else - { - // ignore and reset the error since host pointers produce an error - cudaGetLastError(); - return std::is_same::value; - } -#endif - } - } // namespace ArrayViewHelper -} // namespace internal - - - template inline ArrayView::ArrayView() : starting_element(nullptr) @@ -398,14 +359,7 @@ inline ArrayView::ArrayView( const std::size_t n_elements) : starting_element(starting_element) , n_elements(n_elements) -{ - Assert( - n_elements == 0 || - internal::ArrayViewHelper::is_in_correct_memory_space( - starting_element), - ExcMessage("The memory space indicated by the template parameter " - "and the one derived from the pointer value do not match!")); -} +{} diff --git a/include/deal.II/lac/la_parallel_vector.templates.h b/include/deal.II/lac/la_parallel_vector.templates.h index 179a15d8b3..d4910ca6ef 100644 --- a/include/deal.II/lac/la_parallel_vector.templates.h +++ b/include/deal.II/lac/la_parallel_vector.templates.h @@ -86,9 +86,10 @@ namespace LinearAlgebra template struct la_parallel_vector_templates_functions { - static_assert(std::is_same::value || - std::is_same::value, - "MemorySpace should be Host or CUDA"); + static_assert( + std::is_same::value || + std::is_same::value, + "MemorySpace should be Host or Device"); static void resize_val( @@ -317,24 +318,27 @@ namespace LinearAlgebra }; template - struct la_parallel_vector_templates_functions + struct la_parallel_vector_templates_functions< + Number, + ::dealii::MemorySpace::Device> { using size_type = types::global_dof_index; static void - resize_val(const types::global_dof_index new_alloc_size, - types::global_dof_index & allocated_size, - ::dealii::MemorySpace:: - MemorySpaceData &data, - const MPI_Comm &comm_sm) + resize_val( + const types::global_dof_index new_alloc_size, + types::global_dof_index & allocated_size, + ::dealii::MemorySpace::MemorySpaceData + & data, + const MPI_Comm &comm_sm) { (void)comm_sm; static_assert( std::is_same::value || std::is_same::value, - "Number should be float or double for CUDA memory space"); + "Number should be float or double for Device memory space"); if (new_alloc_size > allocated_size) { @@ -378,26 +382,35 @@ namespace LinearAlgebra // including ghost entries. this is not really efficient right now // because indices are translated twice, once by nth_index_in_set(i) // and once for operator() of tmp_vector - const IndexSet & v_stored = V.get_stored_elements(); - const size_type n_elements = v_stored.n_elements(); - Kokkos::DefaultHostExecutionSpace host_exec; - Kokkos::View indices("indices", n_elements); - Kokkos::parallel_for(Kokkos::RangePolicy(host_exec, 0, n_elements), - KOKKOS_LAMBDA(size_type i) { - indices[i] = communication_pattern->global_to_local( - v_stored.nth_index_in_set(i)); - }); - host_exec.fence(); - + const IndexSet & v_stored = V.get_stored_elements(); + const size_type n_elements = v_stored.n_elements(); + Kokkos::DefaultHostExecutionSpace host_exec; + Kokkos::View indices("indices", + n_elements); + Kokkos::parallel_for( + Kokkos::RangePolicy(host_exec, + 0, + n_elements), + KOKKOS_LAMBDA(size_type i) { + indices[i] = communication_pattern->global_to_local( + v_stored.nth_index_in_set(i)); + }); + host_exec.fence(); + // Move the indices to the device ::dealii::MemorySpace::Device::kokkos_space::execution_space exec; - Kokkos::View indices_dev("indices_dev", n_elements); - Kokkos::deep_copy(exec, indices_dev, indices); - exec.fence(); - + Kokkos::View + indices_dev("indices_dev", n_elements); + Kokkos::deep_copy(exec, indices_dev, indices); + exec.fence(); + // Move the data to the device - Kokkos::View V_dev("V_dev", n_elements); - Kokkos::deep_copy(exec, V_dev, Kokkos::View(V.begin(), n_elements)); + Kokkos::View + V_dev("V_dev", n_elements); + Kokkos::deep_copy( + exec, + V_dev, + Kokkos::View(V.begin(), n_elements)); exec.fence(); // Set the values in tmp_vector @@ -415,17 +428,22 @@ namespace LinearAlgebra // Copy the local elements of tmp_vector to the right place in val IndexSet tmp_index_set = tmp_vector.locally_owned_elements(); const size_type tmp_n_elements = tmp_index_set.n_elements(); - Kokkos::realloc(indices, tmp_n_elements); - Kokkos::parallel_for(Kokkos::RangePolicy(host_exec, 0, n_elements), - KOKKOS_LAMBDA(size_type i) { - indices[i] = locally_owned_elem.index_within_set( - tmp_index_set.nth_index_in_set(i)); + Kokkos::realloc(indices, tmp_n_elements); + Kokkos::parallel_for( + Kokkos::RangePolicy(host_exec, + 0, + n_elements), + KOKKOS_LAMBDA(size_type i) { + indices[i] = locally_owned_elem.index_within_set( + tmp_index_set.nth_index_in_set(i)); }); host_exec.fence(); - Kokkos::realloc(indices_dev, tmp_n_elements); - Kokkos::deep_copy(indices_dev, Kokkos::View(indices.data(), tmp_n_elements)); + Kokkos::realloc(indices_dev, tmp_n_elements); + Kokkos::deep_copy(indices_dev, + Kokkos::View(indices.data(), + tmp_n_elements)); - if (operation == VectorOperation::add) + if (operation == VectorOperation::add) Kokkos::parallel_for( Kokkos::RangePolicy< ::dealii::MemorySpace::Device::kokkos_space::execution_space>( @@ -449,13 +467,13 @@ namespace LinearAlgebra linfty_norm_local(const ::dealii::MemorySpace::MemorySpaceData< Number, ::dealii::MemorySpace::Device> &data, - const unsigned int size, - RealType & result) + const unsigned int size, + RealType & result) { static_assert(std::is_same::value, "RealType should be the same type as Number"); - typename ::dealii::MemorySpace::Device::kokkos_space::execution_space + typename ::dealii::MemorySpace::Device::kokkos_space::execution_space exec; Kokkos::parallel_reduce( Kokkos::RangePolicy< @@ -929,7 +947,7 @@ namespace LinearAlgebra if (data.values_host_buffer.size() > 0) Kokkos::deep_copy(Kokkos::subview(data.values_host_buffer, range), 0); if (data.values.size() > 0) - Kokkos::deep_copy(Kokkos::subview(data.values, range), 0); + Kokkos::deep_copy(Kokkos::subview(data.values, range), 0); vector_is_ghosted = false; } @@ -1192,11 +1210,12 @@ namespace LinearAlgebra // The communication is done on the host, so we need to // move the data back to the device. - auto range = Kokkos::make_pair( - partitioner->locally_owned_size(), - partitioner->locally_owned_size() + partitioner->n_ghost_indices()); + auto range = Kokkos::make_pair(partitioner->locally_owned_size(), + partitioner->locally_owned_size() + + partitioner->n_ghost_indices()); Kokkos::deep_copy(Kokkos::subview(data.values, range), - Kokkos::subview(data.values_host_buffer, range)); + Kokkos::subview(data.values_host_buffer, + range)); Kokkos::resize(data.values_host_buffer, 0); } @@ -1764,7 +1783,6 @@ namespace LinearAlgebra { real_type sum; - dealii::internal::VectorOperations:: functions::norm_2( thread_loop_partitioner, diff --git a/include/deal.II/lac/vector_operations_internal.h b/include/deal.II/lac/vector_operations_internal.h index 21d08f4413..452e378ce0 100644 --- a/include/deal.II/lac/vector_operations_internal.h +++ b/include/deal.II/lac/vector_operations_internal.h @@ -1484,9 +1484,9 @@ namespace internal ::dealii::MemorySpace::MemorySpaceData & /*data*/) { static_assert( - std::is_same::value && + std::is_same::value && std::is_same::value, - "For the CUDA MemorySpace Number and Number2 should be the same type"); + "For the Device MemorySpace Number and Number2 should be the same type"); } static void @@ -2097,7 +2097,6 @@ namespace internal } } -#ifdef DEAL_II_COMPILER_CUDA_AWARE template static void import_elements( @@ -2111,51 +2110,44 @@ namespace internal ::dealii::MemorySpace::Host> &data, std::enable_if_t< - std::is_same::value, + std::is_same::value, int> = 0) { if (operation == VectorOperation::insert) { - cudaError_t cuda_error_code = cudaMemcpy(data.values.data(), - v_data.values.data(), - size * sizeof(Number), - cudaMemcpyDeviceToHost); - AssertCuda(cuda_error_code); + Kokkos::deep_copy( + Kokkos::subview(data.values, + Kokkos::pair(0, size)), + Kokkos::subview(v_data.values, + Kokkos::pair(0, size))); } else { AssertThrow(false, ExcNotImplemented()); } } -#endif }; -#ifdef DEAL_II_COMPILER_CUDA_AWARE template - struct functions + struct functions { - static const int block_size = - ::dealii::LinearAlgebra::CUDAWrappers::kernel::block_size; - static const int chunk_size = - ::dealii::LinearAlgebra::CUDAWrappers::kernel::chunk_size; - static void copy( const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &, const size_type size, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - cudaError_t cuda_error_code = cudaMemcpy(data.values.data(), - v_data.values.data(), - size * sizeof(Number), - cudaMemcpyDeviceToDevice); - AssertCuda(cuda_error_code); + Kokkos::deep_copy( + Kokkos::subview(data.values, + Kokkos::pair(0, size)), + Kokkos::subview(v_data.values, + Kokkos::pair(0, size))); } static void @@ -2163,13 +2155,13 @@ namespace internal const size_type size, const Number s, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::set - <<>>(data.values.data(), s, size); - AssertCudaKernel(); + Kokkos::deep_copy( + Kokkos::subview(data.values, + Kokkos::pair(0, size)), + s); } static void @@ -2177,18 +2169,20 @@ namespace internal const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &, const size_type size, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV - <<>>(data.values.data(), - 1., - v_data.values.data(), - size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(int i) { data.values(i) += v_data.values(i); }); + exec.fence(); } static void @@ -2196,18 +2190,20 @@ namespace internal const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &, const size_type size, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV - <<>>(data.values.data(), - -1., - v_data.values.data(), - size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { data.values(i) -= v_data.values(i); }); + exec.fence(); } static void @@ -2216,13 +2212,18 @@ namespace internal const size_type size, Number a, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_add - <<>>(data.values.data(), a, size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { data.values(i) += a; }); + exec.fence(); } static void @@ -2231,18 +2232,22 @@ namespace internal const size_type size, const Number a, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV - <<>>(data.values.data(), - a, - v_data.values.data(), - size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { + data.values(i) += a * v_data.values(i); + }); + exec.fence(); } static void @@ -2252,22 +2257,24 @@ namespace internal const Number a, const Number b, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, const ::dealii::MemorySpace:: - MemorySpaceData &w_data, + MemorySpaceData &w_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aVbW - <<>>(data.values.data(), - a, - v_data.values.data(), - b, - w_data.values.data(), - size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { + data.values(i) += a * v_data.values(i) + b * w_data.values(i); + }); + exec.fence(); } static void @@ -2276,16 +2283,22 @@ namespace internal const size_type size, const Number x, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd - <<>>( - x, data.values.data(), 1., v_data.values.data(), size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { + data.values(i) = x * data.values(i) + v_data.values(i); + }); + exec.fence(); } static void @@ -2295,16 +2308,22 @@ namespace internal const Number x, const Number a, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd - <<>>( - x, data.values.data(), a, v_data.values.data(), size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { + data.values(i) = x * data.values(i) + a * v_data.values(i); + }); + exec.fence(); } static void @@ -2315,23 +2334,25 @@ namespace internal const Number a, const Number b, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, const ::dealii::MemorySpace:: - MemorySpaceData &w_data, + MemorySpaceData &w_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd - <<>>(x, - data.values.data(), - a, - v_data.values.data(), - b, - w_data.values.data(), - size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { + data.values(i) = + x * data.values(i) + a * v_data.values(i) + b * w_data.values(i); + }); + exec.fence(); } static void @@ -2340,13 +2361,18 @@ namespace internal const size_type size, const Number factor, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_scale - <<>>(data.values.data(), factor, size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { data.values(i) *= factor; }); + exec.fence(); } static void @@ -2354,17 +2380,20 @@ namespace internal const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &, const size_type size, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::scale - <<>>(data.values.data(), - v_data.values.data(), - size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { data.values(i) *= v_data.values(i); }); + exec.fence(); } static void @@ -2373,18 +2402,22 @@ namespace internal const size_type size, const Number a, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::equ - <<>>(data.values.data(), - a, - v_data.values.data(), - size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { + data.values(i) = a * v_data.values(i); + }); + exec.fence(); } static void @@ -2394,63 +2427,50 @@ namespace internal const Number a, const Number b, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, const ::dealii::MemorySpace:: - MemorySpaceData &w_data, + MemorySpaceData &w_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::equ - <<>>(data.values.data(), - a, - v_data.values.data(), - b, - w_data.values.data(), - size); - AssertCudaKernel(); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_for( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i) { + data.values(i) = a * v_data.values(i) + b * w_data.values(i); + }); + exec.fence(); } static Number dot(const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &, const size_type size, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - Number * result_device; - cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number)); - AssertCuda(error_code); - error_code = cudaMemset(result_device, 0, sizeof(Number)); - AssertCuda(error_code); - - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::double_vector_reduction< - Number, - ::dealii::LinearAlgebra::CUDAWrappers::kernel::DotProduct> - <<>>(result_device, - data.values.data(), - v_data.values.data(), - static_cast( - size)); - AssertCudaKernel(); - - // Copy the result back to the host Number result; - error_code = cudaMemcpy(&result, - result_device, - sizeof(Number), - cudaMemcpyDeviceToHost); - AssertCuda(error_code); - // Free the memory on the device - error_code = cudaFree(result_device); - AssertCuda(error_code); - AssertIsFinite(result); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_reduce( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i, Number & update) { + update += data.values(i) * v_data.values(i); + }, + result); + AssertIsFinite(result); return result; } @@ -2460,9 +2480,8 @@ namespace internal & thread_loop_partitioner, const size_type size, real_type & sum, - ::dealii::MemorySpace::MemorySpaceData - &data) + ::dealii::MemorySpace:: + MemorySpaceData &data) { sum = dot(thread_loop_partitioner, size, data, data); } @@ -2472,32 +2491,23 @@ namespace internal const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &, const size_type size, const ::dealii::MemorySpace:: - MemorySpaceData &data) - { - Number * result_device; - cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number)); - AssertCuda(error_code); - error_code = cudaMemset(result_device, 0, sizeof(Number)); - - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction< - Number, - ::dealii::LinearAlgebra::CUDAWrappers::kernel::ElemSum> - <<>>(result_device, - data.values.data(), - size); - - // Copy the result back to the host + MemorySpaceData &data) + { Number result; - error_code = cudaMemcpy(&result, - result_device, - sizeof(Number), - cudaMemcpyDeviceToHost); - AssertCuda(error_code); - // Free the memory on the device - error_code = cudaFree(result_device); - AssertCuda(error_code); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_reduce( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i, Number & update) { + update += data.values(i); + }, + result); + + AssertIsFinite(result); return result; } @@ -2508,44 +2518,44 @@ namespace internal const size_type size, real_type & sum, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - Number * result_device; - cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number)); - AssertCuda(error_code); - error_code = cudaMemset(result_device, 0, sizeof(Number)); - - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction< - Number, - ::dealii::LinearAlgebra::CUDAWrappers::kernel::L1Norm> - <<>>(result_device, - data.values.data(), - size); - - // Copy the result back to the host - error_code = cudaMemcpy(&sum, - result_device, - sizeof(Number), - cudaMemcpyDeviceToHost); - AssertCuda(error_code); - // Free the memory on the device - error_code = cudaFree(result_device); - AssertCuda(error_code); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_reduce( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i, Number & update) { + update += Kokkos::abs(data.values(i)); + }, + sum); } template static void norm_p( const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &, - const size_type, - real_type &, - real_type, + const size_type size, + real_type & sum, + real_type exp, ::dealii::MemorySpace::MemorySpaceData &) + ::dealii::MemorySpace::Device> + &data) { - Assert(false, ExcNotImplemented()); + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_reduce( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i, Number & update) { + update += Kokkos::pow(Kokkos::abs(data.values(i)), exp); + }, + sum); } static Number @@ -2554,33 +2564,29 @@ namespace internal const size_type size, const Number a, const ::dealii::MemorySpace:: - MemorySpaceData &v_data, + MemorySpaceData &v_data, const ::dealii::MemorySpace:: - MemorySpaceData &w_data, + MemorySpaceData &w_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data) { - Number * res_d; - cudaError_t error_code = cudaMalloc(&res_d, sizeof(Number)); - AssertCuda(error_code); - error_code = cudaMemset(res_d, 0, sizeof(Number)); - AssertCuda(error_code); - - const int n_blocks = 1 + size / (chunk_size * block_size); - ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_and_dot - <<>>(res_d, - data.values.data(), - v_data.values.data(), - w_data.values.data(), - a, - size); - Number res; - error_code = - cudaMemcpy(&res, res_d, sizeof(Number), cudaMemcpyDeviceToHost); - AssertCuda(error_code); - error_code = cudaFree(res_d); + + auto exec = typename ::dealii::MemorySpace::Device::kokkos_space:: + execution_space{}; + Kokkos::parallel_reduce( + Kokkos::RangePolicy< + ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec, + 0, + size), + KOKKOS_LAMBDA(size_type i, Number & update) { + data.values(i) += a * v_data.values(i); + update += + data.values(i) * Number(numbers::NumberTraits::conjugate( + w_data.values(i))); + }, + res); return res; } @@ -2595,10 +2601,10 @@ namespace internal const ::dealii::MemorySpace::MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data, std::enable_if_t< - std::is_same::value, + std::is_same::value, int> = 0) { if (operation == VectorOperation::insert) @@ -2625,7 +2631,7 @@ namespace internal const ::dealii::MemorySpace::MemorySpaceData &v_data, ::dealii::MemorySpace::MemorySpaceData + ::dealii::MemorySpace::Device> &data, std::enable_if_t< std::is_same::value, @@ -2633,11 +2639,11 @@ namespace internal { if (operation == VectorOperation::insert) { - cudaError_t cuda_error_code = cudaMemcpy(data.values.data(), - v_data.values.data(), - size * sizeof(Number), - cudaMemcpyHostToDevice); - AssertCuda(cuda_error_code); + Kokkos::deep_copy( + Kokkos::subview(data.values, + Kokkos::pair(0, size)), + Kokkos::subview(v_data.values, + Kokkos::pair(0, size))); } else { @@ -2645,7 +2651,6 @@ namespace internal } } }; -#endif } // namespace VectorOperations } // namespace internal diff --git a/source/lac/CMakeLists.txt b/source/lac/CMakeLists.txt index b12138d840..24f1ad0fff 100644 --- a/source/lac/CMakeLists.txt +++ b/source/lac/CMakeLists.txt @@ -28,6 +28,7 @@ set(_unity_include_src scalapack.cc la_vector.cc la_parallel_vector.cc + la_parallel_vector_cuda.cc la_parallel_block_vector.cc matrix_out.cc precondition_block.cc @@ -169,7 +170,6 @@ if(DEAL_II_WITH_CUDA) cuda_solver_direct.cc cuda_sparse_matrix.cc cuda_vector.cc - la_parallel_vector_cuda.cc read_write_vector_cuda.cc ) endif() diff --git a/source/lac/vector_memory.cc b/source/lac/vector_memory.cc index 98d71dae83..db9ea9dfa2 100644 --- a/source/lac/vector_memory.cc +++ b/source/lac/vector_memory.cc @@ -52,12 +52,12 @@ namespace internal release_all_unused_memory() { #include "vector_memory_release.inst" - dealii::GrowingVectorMemory< - dealii::LinearAlgebra::distributed::Vector>:: - release_unused_memory(); - dealii::GrowingVectorMemory< - dealii::LinearAlgebra::distributed::Vector>:: - release_unused_memory(); + dealii::GrowingVectorMemory>::release_unused_memory(); + dealii::GrowingVectorMemory>::release_unused_memory(); #ifdef DEAL_II_WITH_CUDA release_all_unused_cuda_memory(); #endif -- 2.39.5