cuda_kernels.cc

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.
//
// ------------------------------------------------------------------------
 
#include <deal.II/lac/cuda_kernels.templates.h>
 
DEAL_II_NAMESPACE_OPEN
 
namespace LinearAlgebra
{
  namespace CUDAWrappers
  {
    namespace kernel
    {
      // Explicit instantiation                                              //
 
      template __global__ void
      vec_scale<float>(float *, const float a, const size_type);
      template __global__ void
      vector_bin_op<float, Binop_Addition>(float          *v1,
                                           const float    *v2,
                                           const size_type N);
      template __global__ void
      vector_bin_op<float, Binop_Subtraction>(float          *v1,
                                              const float    *v2,
                                              const size_type N);
      template __global__ void
      masked_vector_bin_op<float, Binop_Addition>(const unsigned int *mask,
                                                  float              *v1,
                                                  const float        *v2,
                                                  const size_type     N);
      template __global__ void
      masked_vector_bin_op<float, Binop_Subtraction>(const unsigned int *mask,
                                                     float              *v1,
                                                     const float        *v2,
                                                     const size_type     N);
      template struct ElemSum<float>;
      template struct L1Norm<float>;
      template struct LInfty<float>;
      template __global__ void
      reduction<float, ElemSum<float>>(float          *result,
                                       const float    *v,
                                       const size_type N);
      template __global__ void
      reduction<float, L1Norm<float>>(float          *result,
                                      const float    *v,
                                      const size_type N);
      template __global__ void
      reduction<float, LInfty<float>>(float          *result,
                                      const float    *v,
                                      const size_type N);
      template struct DotProduct<float>;
      template __global__ void
      double_vector_reduction<float, DotProduct<float>>(float          *result,
                                                        const float    *v1,
                                                        const float    *v2,
                                                        const size_type N);
      template __global__ void
      vec_add<float>(float *val, const float, const size_type N);
      template __global__ void
      add_aV<float>(float          *val,
                    const float     a,
                    const float    *V_val,
                    const size_type N);
      template __global__ void
      add_aVbW<float>(float          *val,
                      const float     a,
                      const float    *V_val,
                      const float     b,
                      const float    *W_val,
                      const size_type N);
      template __global__ void
      sadd<float>(const float     s,
                  float          *val,
                  const float     a,
                  const float    *V_val,
                  const size_type N);
      template __global__ void
      sadd<float>(const float     s,
                  float          *val,
                  const float     a,
                  const float    *V_val,
                  const float     b,
                  const float    *W_val,
                  const size_type N);
      template __global__ void
      scale<float>(float *val, const float *V_val, const size_type N);
      template __global__ void
      equ<float>(float          *val,
                 const float     a,
                 const float    *V_val,
                 const size_type N);
      template __global__ void
      equ<float>(float          *val,
                 const float     a,
                 const float    *V_val,
                 const float     b,
                 const float    *W_val,
                 const size_type N);
      template __global__ void
      add_and_dot<float>(float          *res,
                         float          *v1,
                         const float    *v2,
                         const float    *v3,
                         const float     a,
                         const size_type N);
      template __global__ void
      set<float>(float *val, const float s, const size_type N);
      template __global__ void
      set_permutated<float, size_type>(const size_type *indices,
                                       float           *val,
                                       const float     *v,
                                       const size_type  N);
      template __global__ void
      gather<float, size_type>(float           *val,
                               const size_type *indices,
                               const float     *v,
                               const size_type  N);
      template __global__ void
      add_permutated<float>(const size_type *indices,
                            float           *val,
                            const float     *v,
                            const size_type  N);
 
 
 
      template __global__ void
      vec_scale<double>(double *, const double a, const size_type);
      template __global__ void
      vector_bin_op<double, Binop_Addition>(double         *v1,
                                            const double   *v2,
                                            const size_type N);
      template __global__ void
      vector_bin_op<double, Binop_Subtraction>(double         *v1,
                                               const double   *v2,
                                               const size_type N);
      template __global__ void
      masked_vector_bin_op<double, Binop_Addition>(const unsigned int *mask,
                                                   double             *v1,
                                                   const double       *v2,
                                                   const size_type     N);
      template __global__ void
      masked_vector_bin_op<double, Binop_Subtraction>(const unsigned int *mask,
                                                      double             *v1,
                                                      const double       *v2,
                                                      const size_type     N);
      template struct ElemSum<double>;
      template struct L1Norm<double>;
      template struct LInfty<double>;
      template __global__ void
      reduction<double, ElemSum<double>>(double         *result,
                                         const double   *v,
                                         const size_type N);
      template __global__ void
      reduction<double, L1Norm<double>>(double         *result,
                                        const double   *v,
                                        const size_type N);
      template __global__ void
      reduction<double, LInfty<double>>(double         *result,
                                        const double   *v,
                                        const size_type N);
      template struct DotProduct<double>;
      template __global__ void
      double_vector_reduction<double, DotProduct<double>>(double       *result,
                                                          const double *v1,
                                                          const double *v2,
                                                          const size_type N);
      template __global__ void
      vec_add<double>(double *val, const double, const size_type N);
      template __global__ void
      add_aV<double>(double         *val,
                     const double    a,
                     const double   *V_val,
                     const size_type N);
      template __global__ void
      add_aVbW<double>(double         *val,
                       const double    a,
                       const double   *V_val,
                       const double    b,
                       const double   *W_val,
                       const size_type N);
      template __global__ void
      sadd<double>(const double    s,
                   double         *val,
                   const double    a,
                   const double   *V_val,
                   const size_type N);
      template __global__ void
      sadd<double>(const double    s,
                   double         *val,
                   const double    a,
                   const double   *V_val,
                   const double    b,
                   const double   *W_val,
                   const size_type N);
      template __global__ void
      scale<double>(double *val, const double *V_val, const size_type N);
      template __global__ void
      equ<double>(double         *val,
                  const double    a,
                  const double   *V_val,
                  const size_type N);
      template __global__ void
      equ<double>(double         *val,
                  const double    a,
                  const double   *V_val,
                  const double    b,
                  const double   *W_val,
                  const size_type N);
      template __global__ void
      add_and_dot<double>(double         *res,
                          double         *v1,
                          const double   *v2,
                          const double   *v3,
                          const double    a,
                          const size_type N);
      template __global__ void
      set<double>(double *val, const double s, const size_type N);
      template __global__ void
      set_permutated<double, size_type>(const size_type *indices,
                                        double          *val,
                                        const double    *v,
                                        const size_type  N);
      template __global__ void
      gather<double, size_type>(double          *val,
                                const size_type *indices,
                                const double    *v,
                                const size_type  N);
      template __global__ void
      add_permutated<double>(const size_type *indices,
                             double          *val,
                             const double    *v,
                             const size_type  N);
    } // namespace kernel
  }   // namespace CUDAWrappers
} // namespace LinearAlgebra
 
DEAL_II_NAMESPACE_CLOSE