Reference documentation for deal.II version Git 840703ecf3 2019-10-14 17:15:20 -0400
\(\newcommand{\dealcoloneq}{\mathrel{\vcenter{:}}=}\)
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 size_type = types::global_dof_index;
45 
51  template <typename Number>
52  __global__ void
53  vec_scale(Number *val, const Number a, const size_type N);
54 
55 
56 
62  template <typename Number>
64  {
65  __device__ static inline Number
66  operation(const Number a, const Number b)
67  {
68  return a + b;
69  }
70  };
71 
72  template <typename Number>
73  struct Binop_Addition<std::complex<Number>>
74  {
75  __device__ static inline std::complex<Number>
76  operation(const std::complex<Number> a, const std::complex<Number>)
77  {
78  printf("This function is not implemented for std::complex<Number>!");
79  assert(false);
80  return a;
81  }
82  };
83 
84 
85 
91  template <typename Number>
93  {
94  __device__ static inline Number
95  operation(const Number a, const Number b)
96  {
97  return a - b;
98  }
99  };
100 
101  template <typename Number>
102  struct Binop_Subtraction<std::complex<Number>>
103  {
104  __device__ static inline std::complex<Number>
105  operation(const std::complex<Number> a, const std::complex<Number> b)
106  {
107  printf("This function is not implemented for std::complex<Number>!");
108  assert(false);
109  return a;
110  }
111  };
112 
113 
114 
120  template <typename Number>
121  struct Binop_Max
122  {
123  __device__ static inline Number
124  operation(const Number a, const Number b)
125  {
126  return a > b ? a : b;
127  }
128  };
129 
130  template <typename Number>
131  struct Binop_Max<std::complex<Number>>
132  {
133  __device__ static inline std::complex<Number>
134  operation(const std::complex<Number> a, const std::complex<Number>)
135  {
136  printf("This function is not implemented for std::complex<Number>!");
137  assert(false);
138  return a;
139  }
140  };
141 
142 
143 
149  template <typename Number>
150  struct Binop_Min
151  {
152  __device__ static inline Number
153  operation(const Number a, const Number b)
154  {
155  return a > b ? b : a;
156  }
157  };
158 
159  template <typename Number>
160  struct Binop_Min<std::complex<Number>>
161  {
162  __device__ static inline std::complex<Number>
163  operation(const std::complex<Number> a, const std::complex<Number>)
164  {
165  printf("This function is not implemented for std::complex<Number>!");
166  assert(false);
167  return a;
168  }
169  };
170 
171 
172 
178  template <typename Number, template <typename> class Binop>
179  __global__ void
180  vector_bin_op(Number *v1, const Number *v2, const size_type N);
181 
182 
183 
192  template <typename Number, template <typename> class Binop>
193  __global__ void
194  masked_vector_bin_op(const unsigned int *mask,
195  Number * v1,
196  const Number * v2,
197  const size_type N);
198 
199 
200 
207  template <typename Number>
208  struct ElemSum
209  {
210  __device__ static Number
211  reduction_op(const Number a, const Number b);
212 
213  __device__ static Number
214  atomic_op(Number *dst, const Number a);
215 
216  __device__ static Number
217  element_wise_op(const Number a);
218 
219  __device__ static Number
220  null_value();
221  };
222 
223 
224 
231  template <typename Number>
232  struct L1Norm
233  {
234  __device__ static Number
235  reduction_op(const Number a, const Number b);
236 
237  __device__ static Number
238  atomic_op(Number *dst, const Number a);
239 
240  __device__ static Number
241  element_wise_op(const Number a);
242 
243  __device__ static Number
244  null_value();
245  };
246 
247 
248 
255  template <typename Number>
256  struct LInfty
257  {
258  __device__ static Number
259  reduction_op(const Number a, const Number b);
260 
261  __device__ static Number
262  atomic_op(Number *dst, const Number a);
263 
264  __device__ static Number
265  element_wise_op(const Number a);
266 
267  __device__ static Number
268  null_value();
269  };
270 
271 
272 
278  template <typename Number, typename Operation>
279  __global__ void
280  reduction(Number *result, const Number *v, const size_type N);
281 
282 
283 
290  template <typename Number>
291  struct DotProduct
292  {
293  __device__ static Number
294  binary_op(const Number a, const Number b);
295 
296  __device__ static Number
297  reduction_op(const Number a, const Number b);
298 
299  __device__ static Number
300  atomic_op(Number *dst, const Number a);
301 
302  __device__ static Number
303  null_value();
304  };
305 
306 
307 
314  template <typename Number, typename Operation>
315  __global__ void
316  double_vector_reduction(Number * result,
317  const Number * v1,
318  const Number * v2,
319  const size_type N);
320 
321 
322 
328  template <typename Number>
329  __global__ void
330  vec_add(Number *val, const Number a, const size_type N);
331 
332 
333 
339  template <typename Number>
340  __global__ void
341  add_aV(Number * val,
342  const Number a,
343  const Number * V_val,
344  const size_type N);
345 
346 
347 
354  template <typename Number>
355  __global__ void
356  add_aVbW(Number * val,
357  const Number a,
358  const Number * V_val,
359  const Number b,
360  const Number * W_val,
361  const size_type N);
362 
363 
364 
371  template <typename Number>
372  __global__ void
373  sadd(const Number s,
374  Number * val,
375  const Number a,
376  const Number * V_val,
377  const size_type N);
378 
379 
380 
387  template <typename Number>
388  __global__ void
389  sadd(const Number s,
390  Number * val,
391  const Number a,
392  const Number * V_val,
393  const Number b,
394  const Number * W_val,
395  const size_type N);
396 
397 
398 
405  template <typename Number>
406  __global__ void
407  scale(Number *val, const Number *V_val, const size_type N);
408 
409 
410 
416  template <typename Number>
417  __global__ void
418  equ(Number *val, const Number a, const Number *V_val, const size_type N);
419 
420 
421 
427  template <typename Number>
428  __global__ void
429  equ(Number * val,
430  const Number a,
431  const Number * V_val,
432  const Number b,
433  const Number * W_val,
434  const size_type N);
435 
436 
437 
444  template <typename Number>
445  __global__ void
446  add_and_dot(Number * res,
447  Number * v1,
448  const Number * v2,
449  const Number * v3,
450  const Number a,
451  const size_type N);
452 
453 
454 
460  template <typename Number>
461  __global__ void
462  set(Number *val, const Number s, const size_type N);
463 
464 
471  template <typename Number, typename IndexType>
472  __global__ void
473  set_permutated(const IndexType *indices,
474  Number * val,
475  const Number * v,
476  const IndexType N);
477 
478 
479 
486  template <typename Number, typename IndexType>
487  __global__ void
488  gather(Number * val,
489  const IndexType *indices,
490  const Number * v,
491  const IndexType N);
492 
493 
494 
501  template <typename Number>
502  __global__ void
503  add_permutated(const size_type *indices,
504  Number * val,
505  const Number * v,
506  const size_type N);
507  } // namespace kernel
508  } // namespace CUDAWrappers
509 } // namespace LinearAlgebra
510 
511 DEAL_II_NAMESPACE_CLOSE
512 
513 #endif
514 
515 #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:89
__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)