1 #ifndef PETSC_DEVICE_CUPM_KERNELS_HPP 2 #define PETSC_DEVICE_CUPM_KERNELS_HPP 3 4 #include <petscdevice_cupm.h> 5 6 #if defined(__cplusplus) 7 8 namespace Petsc 9 { 10 11 namespace device 12 { 13 14 namespace cupm 15 { 16 17 namespace kernels 18 { 19 20 namespace util 21 { 22 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 } 29 30 } // namespace util 31 32 } // namespace kernels 33 34 namespace functors 35 { 36 37 template <typename T> 38 class plus_equals { 39 public: 40 using value_type = T; 41 42 PETSC_HOSTDEVICE_DECL constexpr explicit plus_equals(value_type v = value_type{}) noexcept : v_{std::move(v)} { } 43 44 PETSC_NODISCARD PETSC_HOSTDEVICE_INLINE_DECL constexpr value_type operator()(const value_type &val) const noexcept { return val + v_; } 45 46 private: 47 value_type v_; 48 }; 49 50 namespace 51 { 52 53 template <typename T> 54 PETSC_HOSTDEVICE_INLINE_DECL constexpr plus_equals<T> make_plus_equals(const T &v) noexcept 55 { 56 return plus_equals<T>{v}; 57 } 58 59 } // anonymous namespace 60 61 } // namespace functors 62 63 } // namespace cupm 64 65 } // namespace device 66 67 } // namespace Petsc 68 69 #endif // __cplusplus 70 71 #endif // PETSC_DEVICE_CUPM_KERNELS_HPP 72