From 6141476c2c6202063abbab59ccf285f54ee4363e Mon Sep 17 00:00:00 2001 From: Bruno Turcksin Date: Fri, 2 Dec 2022 15:35:32 -0500 Subject: [PATCH] Use DEAL_II_HOST_DEVICE in CUDA matrix-free framework --- .../deal.II/matrix_free/cuda_fe_evaluation.h | 106 ++++++------- .../matrix_free/cuda_hanging_nodes_internal.h | 20 +-- .../deal.II/matrix_free/cuda_matrix_free.h | 39 +++-- .../matrix_free/cuda_matrix_free.templates.h | 36 ++--- .../matrix_free/cuda_tensor_product_kernels.h | 143 +++++++++--------- tests/cuda/coefficient_eval.cc | 4 +- tests/cuda/matrix_free_no_index_initialize.cc | 4 +- tests/cuda/matrix_vector_mf.h | 14 +- 8 files changed, 187 insertions(+), 179 deletions(-) diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h index 39a6214f0f..fd603d9ad9 100644 --- a/include/deal.II/matrix_free/cuda_fe_evaluation.h +++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h @@ -45,14 +45,17 @@ namespace CUDAWrappers * number of points in each space dimensions. */ template - __device__ inline unsigned int + DEAL_II_HOST_DEVICE inline unsigned int compute_index() { - return (dim == 1 ? threadIdx.x % n_points_1d : - dim == 2 ? threadIdx.x % n_points_1d + n_points_1d * threadIdx.y : - threadIdx.x % n_points_1d + - n_points_1d * - (threadIdx.y + n_points_1d * threadIdx.z)); + KOKKOS_IF_ON_DEVICE( + return (dim == 1 ? + threadIdx.x % n_points_1d : + dim == 2 ? + threadIdx.x % n_points_1d + n_points_1d * threadIdx.y : + threadIdx.x % n_points_1d + + n_points_1d * (threadIdx.y + n_points_1d * threadIdx.z));) + KOKKOS_IF_ON_HOST(return 0;) } } // namespace internal @@ -129,7 +132,7 @@ namespace CUDAWrappers /** * Constructor. */ - __device__ + DEAL_II_HOST_DEVICE FEEvaluation(const unsigned int cell_id, const data_type * data, SharedData *shdata); @@ -142,7 +145,7 @@ namespace CUDAWrappers * nodes, so once can see it as a similar function to * AffineConstraints::read_dof_valuess as well. */ - __device__ void + DEAL_II_HOST_DEVICE void read_dof_values(const Number *src); /** @@ -151,7 +154,7 @@ namespace CUDAWrappers * during the write operation. The functionality is hence similar to the * function AffineConstraints::distribute_local_to_global. */ - __device__ void + DEAL_II_HOST_DEVICE void distribute_local_to_global(Number *dst) const; /** @@ -161,7 +164,7 @@ namespace CUDAWrappers * computed. This function needs to be called before the functions * @p get_value() or @p get_gradient() give useful information. */ - __device__ void + DEAL_II_HOST_DEVICE void evaluate(const bool evaluate_val, const bool evaluate_grad); /** @@ -171,49 +174,49 @@ namespace CUDAWrappers * @p integrate_val and @p integrate_grad are used to enable/disable some * of the values or the gradients. */ - __device__ void + DEAL_II_HOST_DEVICE void integrate(const bool integrate_val, const bool integrate_grad); /** * Same as above, except that the quadrature point is computed from thread * id. */ - __device__ value_type + DEAL_II_HOST_DEVICE value_type get_value() const; /** * Same as above, except that the local dof index is computed from the * thread id. */ - __device__ value_type + DEAL_II_HOST_DEVICE value_type get_dof_value() const; /** * Same as above, except that the quadrature point is computed from the * thread id. */ - __device__ void + DEAL_II_HOST_DEVICE void submit_value(const value_type &val_in); /** * Same as above, except that the local dof index is computed from the * thread id. */ - __device__ void + DEAL_II_HOST_DEVICE void submit_dof_value(const value_type &val_in); /** * Same as above, except that the quadrature point is computed from the * thread id. */ - __device__ gradient_type + DEAL_II_HOST_DEVICE gradient_type get_gradient() const; /** * Same as above, except that the quadrature point is computed from the * thread id. */ - __device__ void + DEAL_II_HOST_DEVICE void submit_gradient(const gradient_type &grad_in); // clang-format off @@ -223,13 +226,13 @@ namespace CUDAWrappers * * @p func needs to define * \code - * __device__ void operator()( + * DEAL_II_HOST_DEVICE void operator()( * CUDAWrappers::FEEvaluation *fe_eval) const; * \endcode */ // clang-format on template - __device__ void + DEAL_II_HOST_DEVICE void apply_for_each_quad_point(const Functor &func); private: @@ -258,7 +261,7 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ + DEAL_II_HOST_DEVICE FEEvaluation:: FEEvaluation(const unsigned int cell_id, const data_type * data, @@ -285,7 +288,7 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ void + DEAL_II_HOST_DEVICE void FEEvaluation:: read_dof_values(const Number *src) { @@ -295,9 +298,8 @@ namespace CUDAWrappers const types::global_dof_index src_idx = local_to_global[idx]; // Use the read-only data cache. - values[idx] = __ldg(&src[src_idx]); - - __syncthreads(); + KOKKOS_IF_ON_DEVICE(values[idx] = __ldg(&src[src_idx]); __syncthreads();) + KOKKOS_IF_ON_HOST(values[idx] = src[src_idx];) internal::resolve_hanging_nodes(constraint_mask, values); @@ -310,7 +312,7 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ void + DEAL_II_HOST_DEVICE void FEEvaluation:: distribute_local_to_global(Number *dst) const { @@ -336,7 +338,7 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ void + DEAL_II_HOST_DEVICE void FEEvaluation::evaluate( const bool evaluate_val, const bool evaluate_grad) @@ -354,17 +356,17 @@ namespace CUDAWrappers { evaluator_tensor_product.value_and_gradient_at_quad_pts(values, gradients); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) } else if (evaluate_grad == true) { evaluator_tensor_product.gradient_at_quad_pts(values, gradients); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) } else if (evaluate_val == true) { evaluator_tensor_product.value_at_quad_pts(values); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) } } @@ -375,7 +377,7 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ void + DEAL_II_HOST_DEVICE void FEEvaluation::integrate( const bool integrate_val, const bool integrate_grad) @@ -395,13 +397,13 @@ namespace CUDAWrappers else if (integrate_val == true) { evaluator_tensor_product.integrate_value(values); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) } else if (integrate_grad == true) { evaluator_tensor_product.template integrate_gradient(values, gradients); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) } } @@ -412,11 +414,11 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ typename FEEvaluation::value_type + DEAL_II_HOST_DEVICE typename FEEvaluation::value_type FEEvaluation:: get_value() const { @@ -431,11 +433,11 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ typename FEEvaluation::value_type + DEAL_II_HOST_DEVICE typename FEEvaluation::value_type FEEvaluation:: get_dof_value() const { @@ -450,7 +452,7 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ void + DEAL_II_HOST_DEVICE void FEEvaluation:: submit_value(const value_type &val_in) { @@ -465,7 +467,7 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ void + DEAL_II_HOST_DEVICE void FEEvaluation:: submit_dof_value(const value_type &val_in) { @@ -480,11 +482,11 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ typename FEEvaluation::gradient_type + DEAL_II_HOST_DEVICE typename FEEvaluation::gradient_type FEEvaluation:: get_gradient() const { @@ -514,7 +516,7 @@ namespace CUDAWrappers int n_q_points_1d, int n_components_, typename Number> - __device__ void + DEAL_II_HOST_DEVICE void FEEvaluation:: submit_gradient(const gradient_type &grad_in) { @@ -539,13 +541,13 @@ namespace CUDAWrappers int n_components_, typename Number> template - __device__ void + DEAL_II_HOST_DEVICE void FEEvaluation:: apply_for_each_quad_point(const Functor &func) { func(this); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) } } // namespace CUDAWrappers 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 e15070c511..c354e0b2b7 100644 --- a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h +++ b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h @@ -24,6 +24,8 @@ # include +# include + DEAL_II_NAMESPACE_OPEN namespace CUDAWrappers { @@ -37,7 +39,7 @@ namespace CUDAWrappers // Functions for resolving the hanging node constraints on the GPU // //------------------------------------------------------------------------// template - __device__ inline unsigned int + DEAL_II_HOST_DEVICE inline unsigned int index2(unsigned int i, unsigned int j) { return i + size * j; @@ -46,7 +48,7 @@ namespace CUDAWrappers template - __device__ inline unsigned int + DEAL_II_HOST_DEVICE inline unsigned int index3(unsigned int i, unsigned int j, unsigned int k) { return i + size * j + size * size * k; @@ -58,7 +60,7 @@ namespace CUDAWrappers unsigned int direction, bool transpose, typename Number> - __device__ inline void + DEAL_II_HOST_DEVICE inline void interpolate_boundary_2d( const dealii::internal::MatrixFreeFunctions::ConstraintKinds constraint_mask, @@ -151,11 +153,11 @@ namespace CUDAWrappers // The synchronization is done for all the threads in one block with // each block being assigned to one element. - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) if (constrained_face && constrained_dof) values[index2(x_idx, y_idx)] = t; - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) } @@ -164,7 +166,7 @@ namespace CUDAWrappers unsigned int direction, bool transpose, typename Number> - __device__ inline void + DEAL_II_HOST_DEVICE inline void interpolate_boundary_3d( const dealii::internal::MatrixFreeFunctions::ConstraintKinds constraint_mask, @@ -293,14 +295,14 @@ namespace CUDAWrappers // The synchronization is done for all the threads in one block with // each block being assigned to one element. - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) if ((constrained_face != dealii::internal::MatrixFreeFunctions:: ConstraintKinds::unconstrained) && constrained_dof) values[index3(x_idx, y_idx, z_idx)] = t; - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) } @@ -313,7 +315,7 @@ namespace CUDAWrappers * @cite kronbichler2019multigrid. */ template - __device__ void + DEAL_II_HOST_DEVICE void resolve_hanging_nodes( const dealii::internal::MatrixFreeFunctions::ConstraintKinds constraint_mask, diff --git a/include/deal.II/matrix_free/cuda_matrix_free.h b/include/deal.II/matrix_free/cuda_matrix_free.h index cf2c0a2c4c..fca3ad059b 100644 --- a/include/deal.II/matrix_free/cuda_matrix_free.h +++ b/include/deal.II/matrix_free/cuda_matrix_free.h @@ -38,6 +38,8 @@ # include # include +# include + DEAL_II_NAMESPACE_OPEN @@ -296,7 +298,7 @@ namespace CUDAWrappers * * @p func needs to define * \code - * __device__ void operator()( + * DEAL_II_HOST_DEVICE void operator()( * const unsigned int cell, * const typename CUDAWrappers::MatrixFree::Data *gpu_data, * CUDAWrappers::SharedData * shared_data, @@ -321,7 +323,7 @@ namespace CUDAWrappers * * @p func needs to define * \code - * __device__ void operator()( + * DEAL_II_HOST_DEVICE void operator()( * const unsigned int cell, * const typename CUDAWrappers::MatrixFree::Data *gpu_data); * static const unsigned int n_dofs_1d; @@ -674,7 +676,7 @@ namespace CUDAWrappers /** * Constructor. */ - __device__ + DEAL_II_HOST_DEVICE SharedData(Number *vd, Number *gq[dim]) : values(vd) { @@ -700,7 +702,7 @@ namespace CUDAWrappers // This function determines the number of cells per block, possibly at compile // time (by virtue of being 'constexpr') // TODO this function should be rewritten using meta-programming - __host__ __device__ constexpr unsigned int + DEAL_II_HOST_DEVICE constexpr unsigned int cells_per_block_shmem(int dim, int fe_degree) { /* clang-format off */ @@ -727,14 +729,18 @@ namespace CUDAWrappers * @relates CUDAWrappers::MatrixFree */ template - __device__ inline unsigned int + DEAL_II_HOST_DEVICE inline unsigned int q_point_id_in_cell(const unsigned int n_q_points_1d) { - return ( - dim == 1 ? threadIdx.x % n_q_points_1d : - dim == 2 ? threadIdx.x % n_q_points_1d + n_q_points_1d * threadIdx.y : - threadIdx.x % n_q_points_1d + - n_q_points_1d * (threadIdx.y + n_q_points_1d * threadIdx.z)); + KOKKOS_IF_ON_DEVICE( + return (dim == 1 ? + threadIdx.x % n_q_points_1d : + dim == 2 ? + threadIdx.x % n_q_points_1d + n_q_points_1d * threadIdx.y : + threadIdx.x % n_q_points_1d + + n_q_points_1d * (threadIdx.y + n_q_points_1d * threadIdx.z));) + + KOKKOS_IF_ON_HOST(AssertThrow(false, ExcInternalError()); return 0;) } @@ -746,7 +752,7 @@ namespace CUDAWrappers * @relates CUDAWrappers::MatrixFree */ template - __device__ inline unsigned int + DEAL_II_HOST_DEVICE inline unsigned int local_q_point_id( const unsigned int cell, const typename CUDAWrappers::MatrixFree::Data *data, @@ -765,11 +771,12 @@ namespace CUDAWrappers * @relates CUDAWrappers::MatrixFree */ template - __device__ inline typename CUDAWrappers::MatrixFree::point_type & - get_quadrature_point( - const unsigned int cell, - const typename CUDAWrappers::MatrixFree::Data *data, - const unsigned int n_q_points_1d) + DEAL_II_HOST_DEVICE inline + typename CUDAWrappers::MatrixFree::point_type & + get_quadrature_point( + const unsigned int cell, + const typename CUDAWrappers::MatrixFree::Data *data, + const unsigned int n_q_points_1d) { return *(data->q_points + data->padding_length * cell + q_point_id_in_cell(n_q_points_1d)); 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 1fde44beb0..7ab5b105f8 100644 --- a/include/deal.II/matrix_free/cuda_matrix_free.templates.h +++ b/include/deal.II/matrix_free/cuda_matrix_free.templates.h @@ -75,56 +75,56 @@ namespace CUDAWrappers [data_array_size]; template - __host__ __device__ inline DataArray & - get_global_shape_values(unsigned int i); + DEAL_II_HOST_DEVICE inline DataArray & + get_global_shape_values(unsigned int i); template <> - __host__ __device__ inline DataArray & - get_global_shape_values(unsigned int i) + DEAL_II_HOST_DEVICE inline DataArray & + get_global_shape_values(unsigned int i) { return global_shape_values_d[i]; } template <> - __host__ __device__ inline DataArray & - get_global_shape_values(unsigned int i) + DEAL_II_HOST_DEVICE inline DataArray & + get_global_shape_values(unsigned int i) { return global_shape_values_f[i]; } template - __host__ __device__ inline DataArray & - get_global_shape_gradients(unsigned int i); + DEAL_II_HOST_DEVICE inline DataArray & + get_global_shape_gradients(unsigned int i); template <> - __host__ __device__ inline DataArray & - get_global_shape_gradients(unsigned int i) + DEAL_II_HOST_DEVICE inline DataArray & + get_global_shape_gradients(unsigned int i) { return global_shape_gradients_d[i]; } template <> - __host__ __device__ inline DataArray & - get_global_shape_gradients(unsigned int i) + DEAL_II_HOST_DEVICE inline DataArray & + get_global_shape_gradients(unsigned int i) { return global_shape_gradients_f[i]; } // for collocation methods template - __host__ __device__ inline DataArray & - get_global_co_shape_gradients(unsigned int i); + DEAL_II_HOST_DEVICE inline DataArray & + get_global_co_shape_gradients(unsigned int i); template <> - __host__ __device__ inline DataArray & - get_global_co_shape_gradients(unsigned int i) + DEAL_II_HOST_DEVICE inline DataArray & + get_global_co_shape_gradients(unsigned int i) { return global_co_shape_gradients_d[i]; } template <> - __host__ __device__ inline DataArray & - get_global_co_shape_gradients(unsigned int i) + DEAL_II_HOST_DEVICE inline DataArray & + get_global_co_shape_gradients(unsigned int i) { return global_co_shape_gradients_f[i]; } 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 c423b48dd6..9a93c939d4 100644 --- a/include/deal.II/matrix_free/cuda_tensor_product_kernels.h +++ b/include/deal.II/matrix_free/cuda_tensor_product_kernels.h @@ -83,7 +83,7 @@ namespace CUDAWrappers static constexpr unsigned int n_q_points = Utilities::pow(n_q_points_1d, dim); - __device__ + DEAL_II_HOST_DEVICE EvaluatorTensorProduct(int mf_object_id); /** @@ -91,7 +91,7 @@ namespace CUDAWrappers * points. */ template - __device__ void + DEAL_II_HOST_DEVICE void values(Number shape_values[], const Number *in, Number *out) const; /** @@ -99,40 +99,40 @@ namespace CUDAWrappers * points for a given @p direction. */ template - __device__ void + DEAL_II_HOST_DEVICE void gradients(Number shape_gradients[], const Number *in, Number *out) const; /** * Helper function for values() and gradients(). */ template - __device__ void + DEAL_II_HOST_DEVICE void apply(Number shape_data[], const Number *in, Number *out) const; /** * Evaluate the finite element function at the quadrature points. */ - __device__ void + DEAL_II_HOST_DEVICE void value_at_quad_pts(Number *u); /** * Helper function for integrate(). Integrate the finite element function. */ - __device__ void + DEAL_II_HOST_DEVICE void integrate_value(Number *u); /** * Evaluate the gradients of the finite element function at the quadrature * points. */ - __device__ void + DEAL_II_HOST_DEVICE void gradient_at_quad_pts(const Number *const u, Number *grad_u[dim]); /** * Evaluate the values and the gradients of the finite element function at * the quadrature points. */ - __device__ void + DEAL_II_HOST_DEVICE void value_and_gradient_at_quad_pts(Number *const u, Number *grad_u[dim]); /** @@ -140,14 +140,14 @@ namespace CUDAWrappers * element function. */ template - __device__ void + DEAL_II_HOST_DEVICE void integrate_gradient(Number *u, Number *grad_u[dim]); /** * Helper function for integrate(). Integrate the values and the gradients * of the finite element function. */ - __device__ void + DEAL_II_HOST_DEVICE void integrate_value_and_gradient(Number *u, Number *grad_u[dim]); const int mf_object_id; @@ -156,7 +156,7 @@ namespace CUDAWrappers template - __device__ + DEAL_II_HOST_DEVICE EvaluatorTensorProduct template - __device__ void + DEAL_II_HOST_DEVICE void EvaluatorTensorProduct template - __device__ void + DEAL_II_HOST_DEVICE void EvaluatorTensorProduct template - __device__ void + DEAL_II_HOST_DEVICE void EvaluatorTensorProduct - inline __device__ void + DEAL_II_HOST_DEVICE inline void EvaluatorTensorProduct( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, true, false, true>( get_global_shape_values(mf_object_id), u, u); @@ -278,10 +275,10 @@ namespace CUDAWrappers { values<0, true, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, true, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<2, true, false, true>( get_global_shape_values(mf_object_id), u, u); @@ -298,7 +295,7 @@ namespace CUDAWrappers template - inline __device__ void + DEAL_II_HOST_DEVICE inline void EvaluatorTensorProduct( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, false, false, true>( get_global_shape_values(mf_object_id), u, u); @@ -328,10 +325,10 @@ namespace CUDAWrappers { values<0, false, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, false, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<2, false, false, true>( get_global_shape_values(mf_object_id), u, u); @@ -348,7 +345,7 @@ namespace CUDAWrappers template - inline __device__ void + DEAL_II_HOST_DEVICE inline void EvaluatorTensorProduct( get_global_shape_values(mf_object_id), u, grad_u[1]); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, true, false, true>(get_global_shape_values( mf_object_id), @@ -394,7 +391,7 @@ namespace CUDAWrappers values<0, true, false, false>( get_global_shape_values(mf_object_id), u, grad_u[2]); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, true, false, true>(get_global_shape_values( mf_object_id), @@ -409,7 +406,7 @@ namespace CUDAWrappers grad_u[2], grad_u[2]); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<2, true, false, true>(get_global_shape_values( mf_object_id), @@ -437,7 +434,7 @@ namespace CUDAWrappers template - inline __device__ void + DEAL_II_HOST_DEVICE inline void EvaluatorTensorProduct< evaluate_general, dim, @@ -452,7 +449,7 @@ namespace CUDAWrappers { values<0, true, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) gradients<0, true, false, false>( get_global_co_shape_gradients(mf_object_id), @@ -465,10 +462,10 @@ namespace CUDAWrappers { values<0, true, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, true, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) gradients<0, true, false, false>( get_global_co_shape_gradients(mf_object_id), @@ -485,13 +482,13 @@ namespace CUDAWrappers { values<0, true, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, true, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<2, true, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) gradients<0, true, false, false>( get_global_co_shape_gradients(mf_object_id), @@ -520,7 +517,7 @@ namespace CUDAWrappers template template - inline __device__ void + DEAL_II_HOST_DEVICE inline void EvaluatorTensorProduct( get_global_shape_values(mf_object_id), grad_u[0], u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) gradients<1, false, true, false>( get_global_shape_gradients(mf_object_id), grad_u[1], u); @@ -575,7 +572,7 @@ namespace CUDAWrappers grad_u[2], grad_u[2]); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, false, false, true>(get_global_shape_values( mf_object_id), @@ -590,14 +587,14 @@ namespace CUDAWrappers grad_u[2], grad_u[2]); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<2, false, add, false>( get_global_shape_values(mf_object_id), grad_u[0], u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<2, false, true, false>( get_global_shape_values(mf_object_id), grad_u[1], u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) gradients<2, false, true, false>( get_global_shape_gradients(mf_object_id), grad_u[2], u); @@ -614,7 +611,7 @@ namespace CUDAWrappers template - inline __device__ void + DEAL_II_HOST_DEVICE inline void EvaluatorTensorProduct(mf_object_id), grad_u[0], u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<0, false, false, true>( get_global_shape_values(mf_object_id), u, u); @@ -644,19 +641,19 @@ namespace CUDAWrappers get_global_co_shape_gradients(mf_object_id), grad_u[1], u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) gradients<0, false, true, false>( get_global_co_shape_gradients(mf_object_id), grad_u[0], u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, false, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<0, false, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) break; } @@ -666,27 +663,27 @@ namespace CUDAWrappers get_global_co_shape_gradients(mf_object_id), grad_u[2], u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) gradients<1, false, true, false>( get_global_co_shape_gradients(mf_object_id), grad_u[1], u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) gradients<0, false, true, false>( get_global_co_shape_gradients(mf_object_id), grad_u[0], u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<2, false, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<1, false, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) values<0, false, false, true>( get_global_shape_values(mf_object_id), u, u); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) break; } diff --git a/tests/cuda/coefficient_eval.cc b/tests/cuda/coefficient_eval.cc index 38da36089b..8db177d609 100644 --- a/tests/cuda/coefficient_eval.cc +++ b/tests/cuda/coefficient_eval.cc @@ -35,7 +35,7 @@ class DummyOperator public: DummyOperator() = default; - __device__ void + DEAL_II_HOST_DEVICE void operator()( const unsigned int cell, const typename CUDAWrappers::MatrixFree::Data *gpu_data, @@ -53,7 +53,7 @@ public: template -__device__ void +DEAL_II_HOST_DEVICE void DummyOperator::operator()( const unsigned int cell, const typename CUDAWrappers::MatrixFree::Data *gpu_data, diff --git a/tests/cuda/matrix_free_no_index_initialize.cc b/tests/cuda/matrix_free_no_index_initialize.cc index 0adf975f14..204658a41a 100644 --- a/tests/cuda/matrix_free_no_index_initialize.cc +++ b/tests/cuda/matrix_free_no_index_initialize.cc @@ -38,7 +38,7 @@ public: MatrixFreeTest(const CUDAWrappers::MatrixFree &data_in) : data(data_in){}; - __device__ void + DEAL_II_HOST_DEVICE void operator()( const unsigned int cell, const typename CUDAWrappers::MatrixFree::Data *gpu_data, @@ -51,7 +51,7 @@ public: // set to unit vector fe_eval.submit_dof_value(1.); - __syncthreads(); + KOKKOS_IF_ON_DEVICE(__syncthreads();) fe_eval.evaluate(/*evaluate_values =*/true, /*evaluate_gradients=*/true); #ifndef __APPLE__ diff --git a/tests/cuda/matrix_vector_mf.h b/tests/cuda/matrix_vector_mf.h index baeb5d6e6f..20af64558e 100644 --- a/tests/cuda/matrix_vector_mf.h +++ b/tests/cuda/matrix_vector_mf.h @@ -30,12 +30,12 @@ template class HelmholtzOperatorQuad { public: - __device__ + DEAL_II_HOST_DEVICE HelmholtzOperatorQuad(Number coef) : coef(coef) {} - __device__ void + DEAL_II_HOST_DEVICE void operator()( CUDAWrappers::FEEvaluation *fe_eval) const; @@ -47,7 +47,7 @@ private: template -__device__ void +DEAL_II_HOST_DEVICE void HelmholtzOperatorQuad::operator()( CUDAWrappers::FEEvaluation *fe_eval) const @@ -66,7 +66,7 @@ public: : coef(coefficient) {} - __device__ void + DEAL_II_HOST_DEVICE void operator()( const unsigned int cell, const typename CUDAWrappers::MatrixFree::Data *gpu_data, @@ -86,7 +86,7 @@ public: template -__device__ void +DEAL_II_HOST_DEVICE void HelmholtzOperator::operator()( const unsigned int cell, const typename CUDAWrappers::MatrixFree::Data *gpu_data, @@ -117,7 +117,7 @@ public: : coef(coefficient) {} - __device__ void + DEAL_II_HOST_DEVICE void operator()( const unsigned int cell, const typename CUDAWrappers::MatrixFree::Data *gpu_data); @@ -135,7 +135,7 @@ private: template -__device__ void +DEAL_II_HOST_DEVICE void VaryingCoefficientFunctor::operator()( const unsigned int cell, const typename CUDAWrappers::MatrixFree::Data *gpu_data) -- 2.39.5