Reference documentation for deal.II version 9.5.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\}}\)
Loading...
Searching...
No Matches
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
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 {
42 using ::CUDAWrappers::block_size;
43 using ::CUDAWrappers::chunk_size;
44 using ::CUDAWrappers::warp_size;
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>
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>
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>
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:472
#define DEAL_II_NAMESPACE_CLOSE
Definition config.h:473
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)
STL namespace.
unsigned int global_dof_index
Definition types.h:82
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number >)
static Number operation(const Number a, const Number b)
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number >)
static Number operation(const Number a, const Number b)
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number >)
static Number operation(const Number a, const Number b)
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number >)
static Number operation(const Number a, const Number b)
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)