cuda_kernels.h

Go to the documentation of this file.

// ------------------------------------------------------------------------
//
// SPDX-License-Identifier: LGPL-2.1-or-later
// Copyright (C) 2018 - 2023 by the deal.II authors
//
// This file is part of the deal.II library.
//
// Part of the source code is dual licensed under Apache-2.0 WITH
// LLVM-exception OR LGPL-2.1-or-later. Detailed license information
// governing the source code and code contributions can be found in
// LICENSE.md and CONTRIBUTING.md at the top level directory of deal.II.
//
// ------------------------------------------------------------------------
 
#ifndef dealii_cuda_kernels_h
#define dealii_cuda_kernels_h
 
#include <deal.II/base/config.h>
 
#ifdef DEAL_II_WITH_CUDA
 
 
#  include <deal.II/base/cuda_size.h>
#  include <deal.II/base/types.h>
 
#  include <deal.II/lac/cuda_atomic.h>
 
#  include <assert.h>
 
DEAL_II_NAMESPACE_OPEN
 
namespace LinearAlgebra
{
  namespace CUDAWrappers
  {
    namespace kernel
    {
      using ::CUDAWrappers::block_size;
      using ::CUDAWrappers::chunk_size;
      using ::CUDAWrappers::warp_size;
      using size_type = types::global_dof_index;
 
      template <typename Number>
      __global__ void
      vec_scale(Number *val, const Number a, const size_type N);
 
 
 
      template <typename Number>
      struct Binop_Addition
      {
        __device__ static inline Number
        operation(const Number a, const Number b)
        {
          return a + b;
        }
      };
 
      template <typename Number>
      struct Binop_Addition<std::complex<Number>>
      {
        __device__ static inline std::complex<Number>
        operation(const std::complex<Number> a, const std::complex<Number>)
        {
          printf("This function is not implemented for std::complex<Number>!");
          assert(false);
          return a;
        }
      };
 
 
 
      template <typename Number>
      struct Binop_Subtraction
      {
        __device__ static inline Number
        operation(const Number a, const Number b)
        {
          return a - b;
        }
      };
 
      template <typename Number>
      struct Binop_Subtraction<std::complex<Number>>
      {
        __device__ static inline std::complex<Number>
        operation(const std::complex<Number> a,
                  const std::complex<Number> /*b*/)
        {
          printf("This function is not implemented for std::complex<Number>!");
          assert(false);
          return a;
        }
      };
 
 
 
      template <typename Number>
      struct Binop_Max
      {
        __device__ static inline Number
        operation(const Number a, const Number b)
        {
          return a > b ? a : b;
        }
      };
 
      template <typename Number>
      struct Binop_Max<std::complex<Number>>
      {
        __device__ static inline std::complex<Number>
        operation(const std::complex<Number> a, const std::complex<Number>)
        {
          printf("This function is not implemented for std::complex<Number>!");
          assert(false);
          return a;
        }
      };
 
 
 
      template <typename Number>
      struct Binop_Min
      {
        __device__ static inline Number
        operation(const Number a, const Number b)
        {
          return a > b ? b : a;
        }
      };
 
      template <typename Number>
      struct Binop_Min<std::complex<Number>>
      {
        __device__ static inline std::complex<Number>
        operation(const std::complex<Number> a, const std::complex<Number>)
        {
          printf("This function is not implemented for std::complex<Number>!");
          assert(false);
          return a;
        }
      };
 
 
 
      template <typename Number, template <typename> class Binop>
      __global__ void
      vector_bin_op(Number *v1, const Number *v2, const size_type N);
 
 
 
      template <typename Number, template <typename> class Binop>
      __global__ void
      masked_vector_bin_op(const unsigned int *mask,
                           Number             *v1,
                           const Number       *v2,
                           const size_type     N);
 
 
 
      template <typename Number>
      struct ElemSum
      {
        __device__ static Number
        reduction_op(const Number a, const Number b);
 
        __device__ static Number
        atomic_op(Number *dst, const Number a);
 
        __device__ static Number
        element_wise_op(const Number a);
 
        __device__ static Number
        null_value();
      };
 
 
 
      template <typename Number>
      struct L1Norm
      {
        __device__ static Number
        reduction_op(const Number a, const Number b);
 
        __device__ static Number
        atomic_op(Number *dst, const Number a);
 
        __device__ static Number
        element_wise_op(const Number a);
 
        __device__ static Number
        null_value();
      };
 
 
 
      template <typename Number>
      struct LInfty
      {
        __device__ static Number
        reduction_op(const Number a, const Number b);
 
        __device__ static Number
        atomic_op(Number *dst, const Number a);
 
        __device__ static Number
        element_wise_op(const Number a);
 
        __device__ static Number
        null_value();
      };
 
 
 
      template <typename Number, typename Operation>
      __global__ void
      reduction(Number *result, const Number *v, const size_type N);
 
 
 
      template <typename Number>
      struct DotProduct
      {
        __device__ static Number
        binary_op(const Number a, const Number b);
 
        __device__ static Number
        reduction_op(const Number a, const Number b);
 
        __device__ static Number
        atomic_op(Number *dst, const Number a);
 
        __device__ static Number
        null_value();
      };
 
 
 
      template <typename Number, typename Operation>
      __global__ void
      double_vector_reduction(Number         *result,
                              const Number   *v1,
                              const Number   *v2,
                              const size_type N);
 
 
 
      template <typename Number>
      __global__ void
      vec_add(Number *val, const Number a, const size_type N);
 
 
 
      template <typename Number>
      __global__ void
      add_aV(Number         *val,
             const Number    a,
             const Number   *V_val,
             const size_type N);
 
 
 
      template <typename Number>
      __global__ void
      add_aVbW(Number         *val,
               const Number    a,
               const Number   *V_val,
               const Number    b,
               const Number   *W_val,
               const size_type N);
 
 
 
      template <typename Number>
      __global__ void
      sadd(const Number    s,
           Number         *val,
           const Number    a,
           const Number   *V_val,
           const size_type N);
 
 
 
      template <typename Number>
      __global__ void
      sadd(const Number    s,
           Number         *val,
           const Number    a,
           const Number   *V_val,
           const Number    b,
           const Number   *W_val,
           const size_type N);
 
 
 
      template <typename Number>
      __global__ void
      scale(Number *val, const Number *V_val, const size_type N);
 
 
 
      template <typename Number>
      __global__ void
      equ(Number *val, const Number a, const Number *V_val, const size_type N);
 
 
 
      template <typename Number>
      __global__ void
      equ(Number         *val,
          const Number    a,
          const Number   *V_val,
          const Number    b,
          const Number   *W_val,
          const size_type N);
 
 
 
      template <typename Number>
      __global__ void
      add_and_dot(Number         *res,
                  Number         *v1,
                  const Number   *v2,
                  const Number   *v3,
                  const Number    a,
                  const size_type N);
 
 
 
      template <typename Number>
      __global__ void
      set(Number *val, const Number s, const size_type N);
 
 
      template <typename Number, typename IndexType>
      __global__ void
      set_permutated(const IndexType *indices,
                     Number          *val,
                     const Number    *v,
                     const IndexType  N);
 
 
 
      template <typename Number, typename IndexType>
      __global__ void
      gather(Number          *val,
             const IndexType *indices,
             const Number    *v,
             const IndexType  N);
 
 
 
      template <typename Number>
      __global__ void
      add_permutated(const size_type *indices,
                     Number          *val,
                     const Number    *v,
                     const size_type  N);
    } // namespace kernel
  }   // namespace CUDAWrappers
} // namespace LinearAlgebra
 
DEAL_II_NAMESPACE_CLOSE
 
#endif
 
#endif