Reference documentation for deal.II version 9.4.0
\(\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 - 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
25# include <deal.II/base/types.h>
26
28
29# include <assert.h>
30
32
33namespace 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, 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>
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>
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
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
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
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>
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
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
513
514#endif
515
516#endif
#define DEAL_II_NAMESPACE_OPEN
Definition: config.h:442
#define DEAL_II_NAMESPACE_CLOSE
Definition: config.h:443
const unsigned int v1
Definition: grid_tools.cc:1000
__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
SymmetricTensor< 2, dim, Number > b(const Tensor< 2, dim, Number > &F)
STL namespace.
unsigned int global_dof_index
Definition: types.h:76
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:135
static Number operation(const Number a, const Number b)
Definition: cuda_kernels.h:125
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number >)
Definition: cuda_kernels.h:164
static Number operation(const Number a, const Number b)
Definition: cuda_kernels.h:154
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number > b)
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)