16 #ifndef dealii_cuda_kernels_h 17 #define dealii_cuda_kernels_h 21 #ifdef DEAL_II_COMPILER_CUDA_AWARE 52 template <
typename Number>
63 template <
typename Number>
66 __device__
static inline Number
73 template <
typename Number>
76 __device__
static inline std::complex<Number>
77 operation(
const std::complex<Number> a,
const std::complex<Number>)
79 printf(
"This function is not implemented for std::complex<Number>!");
92 template <
typename Number>
95 __device__
static inline Number
102 template <
typename Number>
105 __device__
static inline std::complex<Number>
106 operation(
const std::complex<Number> a,
const std::complex<Number>
b)
108 printf(
"This function is not implemented for std::complex<Number>!");
121 template <
typename Number>
124 __device__
static inline Number
127 return a > b ? a :
b;
131 template <
typename Number>
134 __device__
static inline std::complex<Number>
135 operation(
const std::complex<Number> a,
const std::complex<Number>)
137 printf(
"This function is not implemented for std::complex<Number>!");
150 template <
typename Number>
153 __device__
static inline Number
156 return a > b ?
b : a;
160 template <
typename Number>
163 __device__
static inline std::complex<Number>
164 operation(
const std::complex<Number> a,
const std::complex<Number>)
166 printf(
"This function is not implemented for std::complex<Number>!");
179 template <
typename Number,
template <
typename>
class Binop>
193 template <
typename Number,
template <
typename>
class Binop>
208 template <
typename Number>
211 __device__
static Number
212 reduction_op(
const Number a,
const Number
b);
214 __device__
static Number
215 atomic_op(Number *dst,
const Number a);
217 __device__
static Number
218 element_wise_op(
const Number a);
220 __device__
static Number
232 template <
typename Number>
235 __device__
static Number
236 reduction_op(
const Number a,
const Number
b);
238 __device__
static Number
239 atomic_op(Number *dst,
const Number a);
241 __device__
static Number
242 element_wise_op(
const Number a);
244 __device__
static Number
256 template <
typename Number>
259 __device__
static Number
260 reduction_op(
const Number a,
const Number
b);
262 __device__
static Number
263 atomic_op(Number *dst,
const Number a);
265 __device__
static Number
266 element_wise_op(
const Number a);
268 __device__
static Number
279 template <
typename Number,
typename Operation>
291 template <
typename Number>
294 __device__
static Number
295 binary_op(
const Number a,
const Number
b);
297 __device__
static Number
298 reduction_op(
const Number a,
const Number b);
300 __device__
static Number
301 atomic_op(Number *dst,
const Number a);
303 __device__
static Number
315 template <
typename Number,
typename Operation>
329 template <
typename Number>
340 template <
typename Number>
344 const Number * V_val,
355 template <
typename Number>
359 const Number * V_val,
361 const Number * W_val,
372 template <
typename Number>
377 const Number * V_val,
388 template <
typename Number>
393 const Number * V_val,
395 const Number * W_val,
406 template <
typename Number>
417 template <
typename Number>
419 equ(Number *val,
const Number a,
const Number *V_val,
const size_type N);
428 template <
typename Number>
432 const Number * V_val,
434 const Number * W_val,
445 template <
typename Number>
461 template <
typename Number>
463 set(Number *val,
const Number s,
const size_type N);
472 template <
typename Number,
typename IndexType>
487 template <
typename Number,
typename IndexType>
490 const IndexType *indices,
502 template <
typename Number>
static Number operation(const Number a, const Number b)
__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)
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number > b)
__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)
static Number operation(const Number a, const Number b)
static Number operation(const Number a, const Number b)
__global__ void masked_vector_bin_op(const unsigned int *mask, Number *v1, const Number *v2, const size_type N)
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 >)
__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)
#define DEAL_II_NAMESPACE_CLOSE
__global__ void equ(Number *val, const Number a, const Number *V_val, const size_type N)
static std::complex< Number > operation(const std::complex< Number > a, const std::complex< Number >)
__global__ void sadd(const Number s, Number *val, const Number a, const Number *V_val, const size_type N)
SymmetricTensor< 2, dim, Number > b(const Tensor< 2, dim, Number > &F)
unsigned int global_dof_index
#define DEAL_II_NAMESPACE_OPEN
__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)