17#ifndef dealii_cuda_tensor_product_kernels_h
18#define dealii_cuda_tensor_product_kernels_h
24#include <deal.II/matrix_free/cuda_matrix_free.templates.h>
26#ifdef DEAL_II_COMPILER_CUDA_AWARE
74 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
81 static constexpr unsigned int dofs_per_cell =
83 static constexpr unsigned int n_q_points =
93 template <
int direction,
bool dof_to_quad,
bool add,
bool in_place>
95 values(Number shape_values[],
const Number *in, Number *out)
const;
101 template <
int direction,
bool dof_to_quad,
bool add,
bool in_place>
103 gradients(Number shape_gradients[],
const Number *in, Number *out)
const;
108 template <
int direction,
bool dof_to_quad,
bool add,
bool in_place>
110 apply(Number shape_data[],
const Number *in, Number *out)
const;
116 value_at_quad_pts(Number *u);
122 integrate_value(Number *u);
129 gradient_at_quad_pts(
const Number *
const u, Number *grad_u[dim]);
136 value_and_gradient_at_quad_pts(Number *
const u, Number *grad_u[dim]);
144 integrate_gradient(Number *u, Number *grad_u[dim]);
151 integrate_value_and_gradient(Number *u, Number *grad_u[dim]);
158 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
164 Number>::EvaluatorTensorProduct(
int object_id)
165 : mf_object_id(object_id)
170 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
171 template <
int direction,
bool dof_to_quad,
bool add,
bool in_place>
177 Number>::values(Number shape_values[],
181 apply<direction, dof_to_quad, add, in_place>(shape_values, in, out);
186 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
187 template <
int direction,
bool dof_to_quad,
bool add,
bool in_place>
193 Number>::gradients(Number shape_gradients[],
197 apply<direction, dof_to_quad, add, in_place>(shape_gradients, in, out);
202 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
203 template <
int direction,
bool dof_to_quad,
bool add,
bool in_place>
209 Number>::apply(Number shape_data[],
213 const unsigned int i = (dim == 1) ? 0 : threadIdx.x % n_q_points_1d;
214 const unsigned int j = (dim == 3) ? threadIdx.y : 0;
215 const unsigned int q = (dim == 1) ? (threadIdx.x % n_q_points_1d) :
216 (dim == 2) ? threadIdx.y :
222 for (
int k = 0; k < n_q_points_1d; ++k)
224 const unsigned int shape_idx =
225 dof_to_quad ? (q + k * n_q_points_1d) : (k + q * n_q_points_1d);
226 const unsigned int source_idx =
227 (direction == 0) ? (k + n_q_points_1d * (i + n_q_points_1d * j)) :
228 (direction == 1) ? (i + n_q_points_1d * (k + n_q_points_1d * j)) :
229 (i + n_q_points_1d * (j + n_q_points_1d * k));
230 t += shape_data[shape_idx] *
231 (in_place ? out[source_idx] : in[source_idx]);
237 const unsigned int destination_idx =
238 (direction == 0) ? (q + n_q_points_1d * (i + n_q_points_1d * j)) :
239 (direction == 1) ? (i + n_q_points_1d * (q + n_q_points_1d * j)) :
240 (i + n_q_points_1d * (j + n_q_points_1d * q));
243 out[destination_idx] += t;
245 out[destination_idx] = t;
250 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
251 inline __device__
void
256 Number>::value_at_quad_pts(Number *u)
262 values<0, true, false, true>(
263 get_global_shape_values<Number>(mf_object_id), u, u);
269 values<0, true, false, true>(
270 get_global_shape_values<Number>(mf_object_id), u, u);
272 values<1, true, false, true>(
273 get_global_shape_values<Number>(mf_object_id), u, u);
279 values<0, true, false, true>(
280 get_global_shape_values<Number>(mf_object_id), u, u);
282 values<1, true, false, true>(
283 get_global_shape_values<Number>(mf_object_id), u, u);
285 values<2, true, false, true>(
286 get_global_shape_values<Number>(mf_object_id), u, u);
300 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
301 inline __device__
void
306 Number>::integrate_value(Number *u)
312 values<0, false, false, true>(
313 get_global_shape_values<Number>(mf_object_id), u, u);
319 values<0, false, false, true>(
320 get_global_shape_values<Number>(mf_object_id), u, u);
322 values<1, false, false, true>(
323 get_global_shape_values<Number>(mf_object_id), u, u);
329 values<0, false, false, true>(
330 get_global_shape_values<Number>(mf_object_id), u, u);
332 values<1, false, false, true>(
333 get_global_shape_values<Number>(mf_object_id), u, u);
335 values<2, false, false, true>(
336 get_global_shape_values<Number>(mf_object_id), u, u);
350 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
351 inline __device__
void
356 Number>::gradient_at_quad_pts(
const Number *
const u,
363 gradients<0, true, false, false>(
364 get_global_shape_gradients<Number>(mf_object_id), u, grad_u[0]);
370 gradients<0, true, false, false>(
371 get_global_shape_gradients<Number>(mf_object_id), u, grad_u[0]);
372 values<0, true, false, false>(
373 get_global_shape_values<Number>(mf_object_id), u, grad_u[1]);
377 values<1, true, false, true>(get_global_shape_values<Number>(
381 gradients<1, true, false, true>(
382 get_global_shape_gradients<Number>(mf_object_id),
390 gradients<0, true, false, false>(
391 get_global_shape_gradients<Number>(mf_object_id), u, grad_u[0]);
392 values<0, true, false, false>(
393 get_global_shape_values<Number>(mf_object_id), u, grad_u[1]);
394 values<0, true, false, false>(
395 get_global_shape_values<Number>(mf_object_id), u, grad_u[2]);
399 values<1, true, false, true>(get_global_shape_values<Number>(
403 gradients<1, true, false, true>(
404 get_global_shape_gradients<Number>(mf_object_id),
407 values<1, true, false, true>(get_global_shape_values<Number>(
414 values<2, true, false, true>(get_global_shape_values<Number>(
418 values<2, true, false, true>(get_global_shape_values<Number>(
422 gradients<2, true, false, true>(
423 get_global_shape_gradients<Number>(mf_object_id),
439 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
440 inline __device__
void
446 Number>::value_and_gradient_at_quad_pts(Number *
const u,
447 Number * grad_u[dim])
453 values<0, true, false, true>(
454 get_global_shape_values<Number>(mf_object_id), u, u);
457 gradients<0, true, false, false>(
458 get_global_co_shape_gradients<Number>(mf_object_id),
466 values<0, true, false, true>(
467 get_global_shape_values<Number>(mf_object_id), u, u);
469 values<1, true, false, true>(
470 get_global_shape_values<Number>(mf_object_id), u, u);
473 gradients<0, true, false, false>(
474 get_global_co_shape_gradients<Number>(mf_object_id),
477 gradients<1, true, false, false>(
478 get_global_co_shape_gradients<Number>(mf_object_id),
486 values<0, true, false, true>(
487 get_global_shape_values<Number>(mf_object_id), u, u);
489 values<1, true, false, true>(
490 get_global_shape_values<Number>(mf_object_id), u, u);
492 values<2, true, false, true>(
493 get_global_shape_values<Number>(mf_object_id), u, u);
496 gradients<0, true, false, false>(
497 get_global_co_shape_gradients<Number>(mf_object_id),
500 gradients<1, true, false, false>(
501 get_global_co_shape_gradients<Number>(mf_object_id),
504 gradients<2, true, false, false>(
505 get_global_co_shape_gradients<Number>(mf_object_id),
521 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
523 inline __device__
void
528 Number>::integrate_gradient(Number *u,
535 gradients<0, false, add, false>(
536 get_global_shape_gradients<Number>(mf_object_id),
544 gradients<0, false, false, true>(
545 get_global_shape_gradients<Number>(mf_object_id),
548 values<0, false, false, true>(get_global_shape_values<Number>(
555 values<1, false, add, false>(
556 get_global_shape_values<Number>(mf_object_id), grad_u[0], u);
558 gradients<1, false, true, false>(
559 get_global_shape_gradients<Number>(mf_object_id), grad_u[1], u);
565 gradients<0, false, false, true>(
566 get_global_shape_gradients<Number>(mf_object_id),
569 values<0, false, false, true>(get_global_shape_values<Number>(
573 values<0, false, false, true>(get_global_shape_values<Number>(
580 values<1, false, false, true>(get_global_shape_values<Number>(
584 gradients<1, false, false, true>(
585 get_global_shape_gradients<Number>(mf_object_id),
588 values<1, false, false, true>(get_global_shape_values<Number>(
595 values<2, false, add, false>(
596 get_global_shape_values<Number>(mf_object_id), grad_u[0], u);
598 values<2, false, true, false>(
599 get_global_shape_values<Number>(mf_object_id), grad_u[1], u);
601 gradients<2, false, true, false>(
602 get_global_shape_gradients<Number>(mf_object_id), grad_u[2], u);
616 template <
int dim,
int fe_degree,
int n_q_po
ints_1d,
typename Number>
617 inline __device__
void
622 Number>::integrate_value_and_gradient(Number *u,
630 gradients<0, false, true, false>(
631 get_global_co_shape_gradients<Number>(mf_object_id),
636 values<0, false, false, true>(
637 get_global_shape_values<Number>(mf_object_id), u, u);
643 gradients<1, false, true, false>(
644 get_global_co_shape_gradients<Number>(mf_object_id),
648 gradients<0, false, true, false>(
649 get_global_co_shape_gradients<Number>(mf_object_id),
654 values<1, false, false, true>(
655 get_global_shape_values<Number>(mf_object_id), u, u);
657 values<0, false, false, true>(
658 get_global_shape_values<Number>(mf_object_id), u, u);
665 gradients<2, false, true, false>(
666 get_global_co_shape_gradients<Number>(mf_object_id),
670 gradients<1, false, true, false>(
671 get_global_co_shape_gradients<Number>(mf_object_id),
675 gradients<0, false, true, false>(
676 get_global_co_shape_gradients<Number>(mf_object_id),
681 values<2, false, false, true>(
682 get_global_shape_values<Number>(mf_object_id), u, u);
684 values<1, false, false, true>(
685 get_global_shape_values<Number>(mf_object_id), u, u);
687 values<0, false, false, true>(
688 get_global_shape_values<Number>(mf_object_id), u, u);
#define DEAL_II_NAMESPACE_OPEN
#define DEAL_II_NAMESPACE_CLOSE
constexpr T pow(const T base, const int iexp)