Reference documentation for deal.II version Git 1d6f7faded 2020-04-08 17:58:29 +0200
\(\newcommand{\vcentcolon}{\mathrel{\mathop{:}}}\) \(\newcommand{\dealcoloneq}{\vcentcolon\mathrel{\mkern-1.2mu}=}\) \(\newcommand{\jump}[1]{\left[\!\left[ #1 \right]\!\right]}\) \(\newcommand{\average}[1]{\left\{\!\left\{ #1 \right\}\!\right\}}\)
cuda_kernels.h
1 // ---------------------------------------------------------------------
2 //
3 // Copyright (C) 2018 - 2019 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_COMPILER_CUDA_AWARE
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 
31 DEAL_II_NAMESPACE_OPEN
32 
33 namespace LinearAlgebra
34 {
35  namespace CUDAWrappers
36  {
40  namespace kernel
41  {
42  using ::CUDAWrappers::block_size;
43  using ::CUDAWrappers::chunk_size;
44  using ::CUDAWrappers::warp_size;
45  using size_type = types::global_dof_index;
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, const std::complex<Number> b)
107  {
108  printf("This function is not implemented for std::complex<Number>!");
109  assert(false);
110  return a;
111  }
112  };
113 
114 
115 
121  template <typename Number>
122  struct Binop_Max
123  {
124  __device__ static inline Number
125  operation(const Number a, const Number b)
126  {
127  return a > b ? a : b;
128  }
129  };
130 
131  template <typename Number>
132  struct Binop_Max<std::complex<Number>>
133  {
134  __device__ static inline std::complex<Number>
135  operation(const std::complex<Number> a, const std::complex<Number>)
136  {
137  printf("This function is not implemented for std::complex<Number>!");
138  assert(false);
139  return a;
140  }
141  };
142 
143 
144 
150  template <typename Number>
151  struct Binop_Min
152  {
153  __device__ static inline Number
154  operation(const Number a, const Number b)
155  {
156  return a > b ? b : a;
157  }
158  };
159 
160  template <typename Number>
161  struct Binop_Min<std::complex<Number>>
162  {
163  __device__ static inline std::complex<Number>
164  operation(const std::complex<Number> a, const std::complex<Number>)
165  {
166  printf("This function is not implemented for std::complex<Number>!");
167  assert(false);
168  return a;
169  }
170  };
171 
172 
173 
179  template <typename Number, template <typename> class Binop>
180  __global__ void
181  vector_bin_op(Number *v1, const Number *v2, const size_type N);
182 
183 
184 
193  template <typename Number, template <typename> class Binop>
194  __global__ void
195  masked_vector_bin_op(const unsigned int *mask,
196  Number * v1,
197  const Number * v2,
198  const size_type N);
199 
200 
201 
208  template <typename Number>
209  struct ElemSum
210  {
211  __device__ static Number
212  reduction_op(const Number a, const Number b);
213 
214  __device__ static Number
215  atomic_op(Number *dst, const Number a);
216 
217  __device__ static Number
218  element_wise_op(const Number a);
219 
220  __device__ static Number
221  null_value();
222  };
223 
224 
225 
232  template <typename Number>
233  struct L1Norm
234  {
235  __device__ static Number
236  reduction_op(const Number a, const Number b);
237 
238  __device__ static Number
239  atomic_op(Number *dst, const Number a);
240 
241  __device__ static Number
242  element_wise_op(const Number a);
243 
244  __device__ static Number
245  null_value();
246  };
247 
248 
249 
256  template <typename Number>
257  struct LInfty
258  {
259  __device__ static Number
260  reduction_op(const Number a, const Number b);
261 
262  __device__ static Number
263  atomic_op(Number *dst, const Number a);
264 
265  __device__ static Number
266  element_wise_op(const Number a);
267 
268  __device__ static Number
269  null_value();
270  };
271 
272 
273 
279  template <typename Number, typename Operation>
280  __global__ void
281  reduction(Number *result, const Number *v, const size_type N);
282 
283 
284 
291  template <typename Number>
292  struct DotProduct
293  {
294  __device__ static Number
295  binary_op(const Number a, const Number b);
296 
297  __device__ static Number
298  reduction_op(const Number a, const Number b);
299 
300  __device__ static Number
301  atomic_op(Number *dst, const Number a);
302 
303  __device__ static Number
304  null_value();
305  };
306 
307 
308 
315  template <typename Number, typename Operation>
316  __global__ void
317  double_vector_reduction(Number * result,
318  const Number * v1,
319  const Number * v2,
320  const size_type N);
321 
322 
323 
329  template <typename Number>
330  __global__ void
331  vec_add(Number *val, const Number a, const size_type N);
332 
333 
334 
340  template <typename Number>
341  __global__ void
342  add_aV(Number * val,
343  const Number a,
344  const Number * V_val,
345  const size_type N);
346 
347 
348 
355  template <typename Number>
356  __global__ void
357  add_aVbW(Number * val,
358  const Number a,
359  const Number * V_val,
360  const Number b,
361  const Number * W_val,
362  const size_type N);
363 
364 
365 
372  template <typename Number>
373  __global__ void
374  sadd(const Number s,
375  Number * val,
376  const Number a,
377  const Number * V_val,
378  const size_type N);
379 
380 
381 
388  template <typename Number>
389  __global__ void
390  sadd(const Number s,
391  Number * val,
392  const Number a,
393  const Number * V_val,
394  const Number b,
395  const Number * W_val,
396  const size_type N);
397 
398 
399 
406  template <typename Number>
407  __global__ void
408  scale(Number *val, const Number *V_val, const size_type N);
409 
410 
411 
417  template <typename Number>
418  __global__ void
419  equ(Number *val, const Number a, const Number *V_val, const size_type N);
420 
421 
422 
428  template <typename Number>
429  __global__ void
430  equ(Number * val,
431  const Number a,
432  const Number * V_val,
433  const Number b,
434  const Number * W_val,
435  const size_type N);
436 
437 
438 
445  template <typename Number>
446  __global__ void
447  add_and_dot(Number * res,
448  Number * v1,
449  const Number * v2,
450  const Number * v3,
451  const Number a,
452  const size_type N);
453 
454 
455 
461  template <typename Number>
462  __global__ void
463  set(Number *val, const Number s, const size_type N);
464 
465 
472  template <typename Number, typename IndexType>
473  __global__ void
474  set_permutated(const IndexType *indices,
475  Number * val,
476  const Number * v,
477  const IndexType N);
478 
479 
480 
487  template <typename Number, typename IndexType>
488  __global__ void
489  gather(Number * val,
490  const IndexType *indices,
491  const Number * v,
492  const IndexType N);
493 
494 
495 
502  template <typename Number>
503  __global__ void
504  add_permutated(const size_type *indices,
505  Number * val,
506  const Number * v,
507  const size_type N);
508  } // namespace kernel
509  } // namespace CUDAWrappers
510 } // namespace LinearAlgebra
511 
512 DEAL_II_NAMESPACE_CLOSE
513 
514 #endif
515 
516 #endif
__global__ void reduction(Number *result, const Number *v, const size_type N)
__global__ void double_vector_reduction(Number *result, const Number *v1, const Number *v2, 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 vec_add(Number *val, const Number a, const size_type N)
__global__ void gather(Number *val, const IndexType *indices, const Number *v, const IndexType N)
__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 scale(Number *val, const Number *V_val, const size_type N)
__global__ void add_aV(Number *val, const Number a, const Number *V_val, const size_type N)
__global__ void masked_vector_bin_op(const unsigned int *mask, Number *v1, const Number *v2, const size_type N)
__global__ void vec_scale(Number *val, const Number a, const size_type N)
__global__ void set_permutated(const IndexType *indices, Number *val, const Number *v, const IndexType N)
__global__ void equ(Number *val, const Number a, 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)
unsigned int global_dof_index
Definition: types.h:77
__global__ void add_permutated(const size_type *indices, Number *val, const Number *v, const size_type N)
__global__ void vector_bin_op(Number *v1, const Number *v2, const size_type N)