Reference documentation for deal.II version GIT 6823c776d5 2023-12-04 20:20:02+00:00
\(\newcommand{\dealvcentcolon}{\mathrel{\mathop{:}}}\) \(\newcommand{\dealcoloneq}{\dealvcentcolon\mathrel{\mkern-1.2mu}=}\) \(\newcommand{\jump}[1]{\left[\!\left[ #1 \right]\!\right]}\) \(\newcommand{\average}[1]{\left\{\!\left\{ #1 \right\}\!\right\}}\)
cuda_kernels.h
Go to the documentation of this file.
1 // ---------------------------------------------------------------------
2 //
3 // Copyright (C) 2018 - 2023 by the deal.II authors
4 //
5 // This file is part of the deal.II library.
6 //
7 // The deal.II library is free software; you can use it, redistribute
8 // it, and/or modify it under the terms of the GNU Lesser General
9 // Public License as published by the Free Software Foundation; either
10 // version 2.1 of the License, or (at your option) any later version.
11 // The full text of the license can be found in the file LICENSE.md at
12 // the top level directory of deal.II.
13 //
14 // ---------------------------------------------------------------------
15 
16 #ifndef dealii_cuda_kernels_h
17 #define dealii_cuda_kernels_h
18 
19 #include <deal.II/base/config.h>
20 
21 #ifdef DEAL_II_WITH_CUDA
22 
23 
24 # include <deal.II/base/cuda_size.h>
25 # include <deal.II/base/types.h>
26 
27 # include <deal.II/lac/cuda_atomic.h>
28 
29 # include <assert.h>
30 
32 
33 namespace LinearAlgebra
34 {
35  namespace CUDAWrappers
36  {
40  namespace kernel
41  {
46 
52  template <typename Number>
53  __global__ void
54  vec_scale(Number *val, const Number a, const size_type N);
55 
56 
57 
63  template <typename Number>
65  {
66  __device__ static inline Number
67  operation(const Number a, const Number b)
68  {
69  return a + b;
70  }
71  };
72 
73  template <typename Number>
74  struct Binop_Addition<std::complex<Number>>
75  {
76  __device__ static inline std::complex<Number>
77  operation(const std::complex<Number> a, const std::complex<Number>)
78  {
79  printf("This function is not implemented for std::complex<Number>!");
80  assert(false);
81  return a;
82  }
83  };
84 
85 
86 
92  template <typename Number>
94  {
95  __device__ static inline Number
96  operation(const Number a, const Number b)
97  {
98  return a - b;
99  }
100  };
101 
102  template <typename Number>
103  struct Binop_Subtraction<std::complex<Number>>
104  {
105  __device__ static inline std::complex<Number>
106  operation(const std::complex<Number> a,
107  const std::complex<Number> /*b*/)
108  {
109  printf("This function is not implemented for std::complex<Number>!");
110  assert(false);
111  return a;
112  }
113  };
114 
115 
116 
122  template <typename Number>
123  struct Binop_Max
124  {
125  __device__ static inline Number
126  operation(const Number a, const Number b)
127  {
128  return a > b ? a : b;
129  }
130  };
131 
132  template <typename Number>
133  struct Binop_Max<std::complex<Number>>
134  {
135  __device__ static inline std::complex<Number>
136  operation(const std::complex<Number> a, const std::complex<Number>)
137  {
138  printf("This function is not implemented for std::complex<Number>!");
139  assert(false);
140  return a;
141  }
142  };
143 
144 
145 
151  template <typename Number>
152  struct Binop_Min
153  {
154  __device__ static inline Number
155  operation(const Number a, const Number b)
156  {
157  return a > b ? b : a;
158  }
159  };
160 
161  template <typename Number>
162  struct Binop_Min<std::complex<Number>>
163  {
164  __device__ static inline std::complex<Number>
165  operation(const std::complex<Number> a, const std::complex<Number>)
166  {
167  printf("This function is not implemented for std::complex<Number>!");
168  assert(false);
169  return a;
170  }
171  };
172 
173 
174 
180  template <typename Number, template <typename> class Binop>
181  __global__ void
182  vector_bin_op(Number *v1, const Number *v2, const size_type N);
183 
184 
185 
194  template <typename Number, template <typename> class Binop>
195  __global__ void
196  masked_vector_bin_op(const unsigned int *mask,
197  Number *v1,
198  const Number *v2,
199  const size_type N);
200 
201 
202 
209  template <typename Number>
210  struct ElemSum
211  {
212  __device__ static Number
213  reduction_op(const Number a, const Number b);
214 
215  __device__ static Number
216  atomic_op(Number *dst, const Number a);
217 
218  __device__ static Number
219  element_wise_op(const Number a);
220 
221  __device__ static Number
223  };
224 
225 
226 
233  template <typename Number>
234  struct L1Norm
235  {
236  __device__ static Number
237  reduction_op(const Number a, const Number b);
238 
239  __device__ static Number
240  atomic_op(Number *dst, const Number a);
241 
242  __device__ static Number
243  element_wise_op(const Number a);
244 
245  __device__ static Number
247  };
248 
249 
250 
257  template <typename Number>
258  struct LInfty
259  {
260  __device__ static Number
261  reduction_op(const Number a, const Number b);
262 
263  __device__ static Number
264  atomic_op(Number *dst, const Number a);
265 
266  __device__ static Number
267  element_wise_op(const Number a);
268 
269  __device__ static Number
271  };
272 
273 
274 
280  template <typename Number, typename Operation>
281  __global__ void
282  reduction(Number *result, const Number *v, const size_type N);
283 
284 
285 
292  template <typename Number>
293  struct DotProduct
294  {
295  __device__ static Number
296  binary_op(const Number a, const Number b);
297 
298  __device__ static Number
299  reduction_op(const Number a, const Number b);
300 
301  __device__ static Number
302  atomic_op(Number *dst, const Number a);
303 
304  __device__ static Number
306  };
307 
308 
309 
316  template <typename Number, typename Operation>
317  __global__ void
318  double_vector_reduction(Number *result,
319  const Number *v1,
320  const Number *v2,
321  const size_type N);
322 
323 
324 
330  template <typename Number>
331  __global__ void
332  vec_add(Number *val, const Number a, const size_type N);
333 
334 
335 
341  template <typename Number>
342  __global__ void
343  add_aV(Number *val,
344  const Number a,
345  const Number *V_val,
346  const size_type N);
347 
348 
349 
356  template <typename Number>
357  __global__ void
358  add_aVbW(Number *val,
359  const Number a,
360  const Number *V_val,
361  const Number b,
362  const Number *W_val,
363  const size_type N);
364 
365 
366 
373  template <typename Number>
374  __global__ void
375  sadd(const Number s,
376  Number *val,
377  const Number a,
378  const Number *V_val,
379  const size_type N);
380 
381 
382 
389  template <typename Number>
390  __global__ void
391  sadd(const Number s,
392  Number *val,
393  const Number a,
394  const Number *V_val,
395  const Number b,
396  const Number *W_val,
397  const size_type N);
398 
399 
400 
407  template <typename Number>
408  __global__ void
409  scale(Number *val, const Number *V_val, const size_type N);
410 
411 
412 
418  template <typename Number>
419  __global__ void
420  equ(Number *val, const Number a, const Number *V_val, const size_type N);
421 
422 
423 
429  template <typename Number>
430  __global__ void
431  equ(Number *val,
432  const Number a,
433  const Number *V_val,
434  const Number b,
435  const Number *W_val,
436  const size_type N);
437 
438 
439 
446  template <typename Number>
447  __global__ void
448  add_and_dot(Number *res,
449  Number *v1,
450  const Number *v2,
451  const Number *v3,
452  const Number a,
453  const size_type N);
454 
455 
456 
462  template <typename Number>
463  __global__ void
464  set(Number *val, const Number s, const size_type N);
465 
466 
473  template <typename Number, typename IndexType>
474  __global__ void
475  set_permutated(const IndexType *indices,
476  Number *val,
477  const Number *v,
478  const IndexType N);
479 
480 
481 
488  template <typename Number, typename IndexType>
489  __global__ void
490  gather(Number *val,
491  const IndexType *indices,
492  const Number *v,
493  const IndexType N);
494 
495 
496 
503  template <typename Number>
504  __global__ void
505  add_permutated(const size_type *indices,
506  Number *val,
507  const Number *v,
508  const size_type N);
509  } // namespace kernel
510  } // namespace CUDAWrappers
511 } // namespace LinearAlgebra
512 
514 
515 #endif
516 
517 #endif
#define DEAL_II_NAMESPACE_OPEN
Definition: config.h:477
#define DEAL_II_NAMESPACE_CLOSE
Definition: config.h:478
const unsigned int v1
__global__ void add_aVbW(Number *val, const Number a, const Number *V_val, const Number b, const Number *W_val, const size_type N)
__global__ void set_permutated(const IndexType *indices, Number *val, const Number *v, const IndexType N)
__global__ void add_aV(Number *val, const Number a, const Number *V_val, const size_type N)
__global__ void equ(Number *val, const Number a, const Number *V_val, const size_type N)
__global__ void gather(Number *val, const IndexType *indices, const Number *v, const IndexType N)
__global__ void masked_vector_bin_op(const unsigned int *mask, Number *v1, const Number *v2, const size_type N)
__global__ void vec_add(Number *val, const Number a, const size_type N)
__global__ void double_vector_reduction(Number *result, const Number *v1, const Number *v2, const size_type N)
__global__ void vec_scale(Number *val, const Number a, const size_type N)
__global__ void reduction(Number *result, const Number *v, const size_type N)
__global__ void add_and_dot(Number *res, Number *v1, const Number *v2, const Number *v3, const Number a, const size_type N)
__global__ void set(Number *val, const Number s, const size_type N)
__global__ void vector_bin_op(Number *v1, const Number *v2, const size_type N)
__global__ void scale(Number *val, const Number *V_val, const size_type N)
__global__ void sadd(const Number s, Number *val, const Number a, const Number *V_val, const size_type N)
__global__ void add_permutated(const size_type *indices, Number *val, const Number *v, const size_type N)
constexpr int warp_size
Definition: cuda_size.h:40
constexpr int chunk_size
Definition: cuda_size.h:35
constexpr int block_size
Definition: cuda_size.h:29
static const char N
types::global_dof_index size_type
Definition: cuda_kernels.h:45
SymmetricTensor< 2, dim, Number > b(const Tensor< 2, dim, Number > &F)
unsigned int global_dof_index
Definition: types.h:82
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number >)
Definition: cuda_kernels.h:77
static Number operation(const Number a, const Number b)
Definition: cuda_kernels.h:67
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number >)
Definition: cuda_kernels.h:136
static Number operation(const Number a, const Number b)
Definition: cuda_kernels.h:126
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number >)
Definition: cuda_kernels.h:165
static Number operation(const Number a, const Number b)
Definition: cuda_kernels.h:155
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number >)
Definition: cuda_kernels.h:106
static Number operation(const Number a, const Number b)
Definition: cuda_kernels.h:96
static Number reduction_op(const Number a, const Number b)
static Number atomic_op(Number *dst, const Number a)
static Number binary_op(const Number a, const Number b)
static Number element_wise_op(const Number a)
static Number atomic_op(Number *dst, const Number a)
static Number reduction_op(const Number a, const Number b)
static Number atomic_op(Number *dst, const Number a)
static Number reduction_op(const Number a, const Number b)
static Number element_wise_op(const Number a)
static Number reduction_op(const Number a, const Number b)
static Number atomic_op(Number *dst, const Number a)
static Number element_wise_op(const Number a)