From: Bruno Turcksin Date: Thu, 8 Aug 2019 21:25:52 +0000 (+0000) Subject: Allow to use atomic operations instead of graph coloring in X-Git-Tag: v9.2.0-rc1~1262^2~2 X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=ce2026ae674a88b8b16830226aafcf360982dad1;p=dealii.git Allow to use atomic operations instead of graph coloring in CUDAWrappers::MatrixFree --- diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h index 755b6fdcda..2bbc299a5c 100644 --- a/include/deal.II/matrix_free/cuda_fe_evaluation.h +++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h @@ -23,6 +23,7 @@ # include # include +# include # include # include @@ -199,6 +200,8 @@ namespace CUDAWrappers const unsigned int constraint_mask; + const bool use_coloring; + Number *inv_jac; Number *JxW; @@ -222,6 +225,7 @@ namespace CUDAWrappers : n_cells(data->n_cells) , padding_length(data->padding_length) , constraint_mask(data->constraint_mask[cell_id]) + , use_coloring(data->use_coloring) , values(shdata->values) { local_to_global = data->local_to_global + padding_length * cell_id; @@ -282,7 +286,11 @@ namespace CUDAWrappers (dim > 2 ? threadIdx.z : 0) * n_q_points_1d * n_q_points_1d; const types::global_dof_index destination_idx = local_to_global[idx]; - dst[destination_idx] += values[idx]; + if (use_coloring) + dst[destination_idx] += values[idx]; + else + LinearAlgebra::CUDAWrappers::atomicAdd_wrapper(&dst[destination_idx], + values[idx]); } diff --git a/include/deal.II/matrix_free/cuda_matrix_free.h b/include/deal.II/matrix_free/cuda_matrix_free.h index 2015252c16..d1e07eeaa0 100644 --- a/include/deal.II/matrix_free/cuda_matrix_free.h +++ b/include/deal.II/matrix_free/cuda_matrix_free.h @@ -91,14 +91,19 @@ namespace CUDAWrappers parallel_over_elem }; + /** + * Standardized data struct to pipe additional data to MatrixFree. + */ struct AdditionalData { AdditionalData( const ParallelizationScheme parallelization_scheme = parallel_in_elem, const UpdateFlags mapping_update_flags = update_gradients | - update_JxW_values) + update_JxW_values, + const bool use_coloring = false) : parallelization_scheme(parallelization_scheme) , mapping_update_flags(mapping_update_flags) + , use_coloring(use_coloring) {} /** @@ -120,6 +125,13 @@ namespace CUDAWrappers * must be specified by this field. */ UpdateFlags mapping_update_flags; + + /** + * If true, use graph coloring. Otherwise, use atomic operations. Graph + * coloring ensures bitwise reproducibility but is slower on Pascal and + * newer architectures. + */ + bool use_coloring; }; /** @@ -136,6 +148,7 @@ namespace CUDAWrappers unsigned int padding_length; unsigned int row_start; unsigned int * constraint_mask; + bool use_coloring; }; /** @@ -377,6 +390,13 @@ namespace CUDAWrappers */ ParallelizationScheme parallelization_scheme; + /** + * If true, use graph coloring. Otherwise, use atomic operations. Graph + * coloring ensures bitwise reproducibility but is slower on Pascal and + * newer architectures. + */ + bool use_coloring; + /** * Total number of degrees of freedom. */ diff --git a/include/deal.II/matrix_free/cuda_matrix_free.templates.h b/include/deal.II/matrix_free/cuda_matrix_free.templates.h index 2be980af56..561f5b1861 100644 --- a/include/deal.II/matrix_free/cuda_matrix_free.templates.h +++ b/include/deal.II/matrix_free/cuda_matrix_free.templates.h @@ -588,6 +588,7 @@ namespace CUDAWrappers data_copy.n_cells = n_cells[color]; data_copy.padding_length = padding_length; data_copy.row_start = row_start[color]; + data_copy.use_coloring = use_coloring; return data_copy; } @@ -803,6 +804,7 @@ namespace CUDAWrappers AssertThrow(false, ExcMessage("Invalid parallelization scheme.")); this->parallelization_scheme = additional_data.parallelization_scheme; + this->use_coloring = additional_data.use_coloring; // TODO: only free if we actually need arrays of different length free(); @@ -861,13 +863,27 @@ namespace CUDAWrappers CellFilter begin(IteratorFilters::LocallyOwnedCell(), dof_handler.begin_active()); CellFilter end(IteratorFilters::LocallyOwnedCell(), dof_handler.end()); - const auto fun = [&](const CellFilter &filter) { - return internal::get_conflict_indices(filter, constraints); - }; - std::vector> graph; + if (begin != end) - graph = GraphColoring::make_graph_coloring(begin, end, fun); + { + if (additional_data.use_coloring) + { + const auto fun = [&](const CellFilter &filter) { + return internal::get_conflict_indices(filter, + constraints); + }; + graph = GraphColoring::make_graph_coloring(begin, end, fun); + } + else + { + // If we are not using coloring, all the cells belong to the same + // color. + graph.resize(1, std::vector()); + for (auto cell = begin; cell != end; ++cell) + graph[0].emplace_back(cell); + } + } n_colors = graph.size(); helper.setup_color_arrays(n_colors);