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