Actual source code: kernels.hpp

  1: #ifndef PETSC_DEVICE_CUPM_KERNELS_HPP
  2: #define PETSC_DEVICE_CUPM_KERNELS_HPP

  4: #include <petscdevice_cupm.h>

  6: #if defined(__cplusplus)

  8: namespace Petsc
  9: {

 11: namespace device
 12: {

 14: namespace cupm
 15: {

 17: namespace kernels
 18: {

 20: namespace util
 21: {

 23: template <typename SizeType, typename T>
 24: PETSC_DEVICE_INLINE_DECL static void grid_stride_1D(const SizeType size, T &&func) noexcept
 25: {
 26:   for (SizeType i = blockIdx.x * blockDim.x + threadIdx.x; i < size; i += blockDim.x * gridDim.x) func(i);
 27:   return;
 28: }

 30: } // namespace util

 32: } // namespace kernels

 34: namespace functors
 35: {

 37: template <typename T>
 38: class plus_equals {
 39: public:
 40:   using value_type = T;

 42:   PETSC_HOSTDEVICE_DECL constexpr explicit plus_equals(value_type v = value_type{}) noexcept : v_{std::move(v)} { }

 44:   PETSC_NODISCARD PETSC_HOSTDEVICE_INLINE_DECL constexpr value_type operator()(const value_type &val) const noexcept { return val + v_; }

 46: private:
 47:   value_type v_;
 48: };

 50: template <typename T>
 51: class times_equals {
 52: public:
 53:   using value_type = T;

 55:   PETSC_HOSTDEVICE_DECL constexpr explicit times_equals(value_type v = value_type{}) noexcept : v_{std::move(v)} { }

 57:   PETSC_NODISCARD PETSC_HOSTDEVICE_INLINE_DECL constexpr value_type operator()(const value_type &val) const noexcept { return val * v_; }

 59: private:
 60:   value_type v_;
 61: };

 63: template <typename T>
 64: class axpy {
 65: public:
 66:   using value_type = T;

 68:   PETSC_HOSTDEVICE_DECL constexpr explicit axpy(value_type v = value_type{}) noexcept : v_{std::move(v)} { }

 70:   PETSC_NODISCARD PETSC_HOSTDEVICE_INLINE_DECL constexpr value_type operator()(const value_type &x, const value_type &y) const noexcept { return v_ * x + y; }

 72: private:
 73:   value_type v_;
 74: };

 76: namespace
 77: {

 79: template <typename T>
 80: PETSC_HOSTDEVICE_INLINE_DECL constexpr plus_equals<T> make_plus_equals(const T &v) noexcept
 81: {
 82:   return plus_equals<T>{v};
 83: }

 85: template <typename T>
 86: PETSC_HOSTDEVICE_INLINE_DECL constexpr times_equals<T> make_times_equals(const T &v) noexcept
 87: {
 88:   return times_equals<T>{v};
 89: }

 91: template <typename T>
 92: PETSC_HOSTDEVICE_INLINE_DECL constexpr axpy<T> make_axpy(const T &v) noexcept
 93: {
 94:   return axpy<T>{v};
 95: }

 97: } // anonymous namespace

 99: } // namespace functors

101: } // namespace cupm

103: } // namespace device

105: } // namespace Petsc

107: #endif // __cplusplus

109: #endif // PETSC_DEVICE_CUPM_KERNELS_HPP