Reference documentation for deal.II version GIT relicensing-480-geae235577e 2024-04-25 01:00: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\}}\)
Loading...
Searching...
No Matches
cuda_kernels.h
Go to the documentation of this file.
1// ------------------------------------------------------------------------
2//
3// SPDX-License-Identifier: LGPL-2.1-or-later
4// Copyright (C) 2018 - 2023 by the deal.II authors
5//
6// This file is part of the deal.II library.
7//
8// Part of the source code is dual licensed under Apache-2.0 WITH
9// LLVM-exception OR LGPL-2.1-or-later. Detailed license information
10// governing the source code and code contributions can be found in
11// LICENSE.md and CONTRIBUTING.md at the top level directory of deal.II.
12//
13// ------------------------------------------------------------------------
14
15#ifndef dealii_cuda_kernels_h
16#define dealii_cuda_kernels_h
17
18#include <deal.II/base/config.h>
19
20#ifdef DEAL_II_WITH_CUDA
21
22
24# include <deal.II/base/types.h>
25
27
28# include <assert.h>
29
31
32namespace LinearAlgebra
33{
34 namespace CUDAWrappers
35 {
39 namespace kernel
40 {
41 using ::CUDAWrappers::block_size;
42 using ::CUDAWrappers::chunk_size;
43 using ::CUDAWrappers::warp_size;
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,
106 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
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:502
#define DEAL_II_NAMESPACE_CLOSE
Definition config.h:503
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:81
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)