16d54fb17SJacob Faibussowitsch #ifndef PETSC_DEVICE_CUPM_KERNELS_HPP 26d54fb17SJacob Faibussowitsch #define PETSC_DEVICE_CUPM_KERNELS_HPP 36d54fb17SJacob Faibussowitsch 46d54fb17SJacob Faibussowitsch #include <petscdevice_cupm.h> 56d54fb17SJacob Faibussowitsch 66d54fb17SJacob Faibussowitsch #if defined(__cplusplus) 76d54fb17SJacob Faibussowitsch 86d54fb17SJacob Faibussowitsch namespace Petsc 96d54fb17SJacob Faibussowitsch { 106d54fb17SJacob Faibussowitsch 116d54fb17SJacob Faibussowitsch namespace device 126d54fb17SJacob Faibussowitsch { 136d54fb17SJacob Faibussowitsch 146d54fb17SJacob Faibussowitsch namespace cupm 156d54fb17SJacob Faibussowitsch { 166d54fb17SJacob Faibussowitsch 176d54fb17SJacob Faibussowitsch namespace kernels 186d54fb17SJacob Faibussowitsch { 196d54fb17SJacob Faibussowitsch 206d54fb17SJacob Faibussowitsch namespace util 216d54fb17SJacob Faibussowitsch { 226d54fb17SJacob Faibussowitsch 236d54fb17SJacob Faibussowitsch template <typename SizeType, typename T> 246d54fb17SJacob Faibussowitsch PETSC_DEVICE_INLINE_DECL static void grid_stride_1D(const SizeType size, T &&func) noexcept 256d54fb17SJacob Faibussowitsch { 266d54fb17SJacob Faibussowitsch for (SizeType i = blockIdx.x * blockDim.x + threadIdx.x; i < size; i += blockDim.x * gridDim.x) func(i); 276d54fb17SJacob Faibussowitsch return; 286d54fb17SJacob Faibussowitsch } 296d54fb17SJacob Faibussowitsch 306d54fb17SJacob Faibussowitsch } // namespace util 316d54fb17SJacob Faibussowitsch 326d54fb17SJacob Faibussowitsch } // namespace kernels 336d54fb17SJacob Faibussowitsch 34*90585354SJacob Faibussowitsch namespace functors 35*90585354SJacob Faibussowitsch { 36*90585354SJacob Faibussowitsch 37*90585354SJacob Faibussowitsch template <typename T> 38*90585354SJacob Faibussowitsch class plus_equals { 39*90585354SJacob Faibussowitsch public: 40*90585354SJacob Faibussowitsch using value_type = T; 41*90585354SJacob Faibussowitsch 42*90585354SJacob Faibussowitsch PETSC_HOSTDEVICE_DECL constexpr explicit plus_equals(value_type v = value_type{}) noexcept : v_{std::move(v)} { } 43*90585354SJacob Faibussowitsch 44*90585354SJacob Faibussowitsch PETSC_NODISCARD PETSC_HOSTDEVICE_INLINE_DECL constexpr value_type operator()(const value_type &val) const noexcept { return val + v_; } 45*90585354SJacob Faibussowitsch 46*90585354SJacob Faibussowitsch private: 47*90585354SJacob Faibussowitsch value_type v_; 48*90585354SJacob Faibussowitsch }; 49*90585354SJacob Faibussowitsch 50*90585354SJacob Faibussowitsch namespace 51*90585354SJacob Faibussowitsch { 52*90585354SJacob Faibussowitsch 53*90585354SJacob Faibussowitsch template <typename T> 54*90585354SJacob Faibussowitsch PETSC_HOSTDEVICE_INLINE_DECL constexpr plus_equals<T> make_plus_equals(const T &v) noexcept 55*90585354SJacob Faibussowitsch { 56*90585354SJacob Faibussowitsch return plus_equals<T>{v}; 57*90585354SJacob Faibussowitsch } 58*90585354SJacob Faibussowitsch 59*90585354SJacob Faibussowitsch } // anonymous namespace 60*90585354SJacob Faibussowitsch 61*90585354SJacob Faibussowitsch } // namespace functors 62*90585354SJacob Faibussowitsch 636d54fb17SJacob Faibussowitsch } // namespace cupm 646d54fb17SJacob Faibussowitsch 656d54fb17SJacob Faibussowitsch } // namespace device 666d54fb17SJacob Faibussowitsch 676d54fb17SJacob Faibussowitsch } // namespace Petsc 686d54fb17SJacob Faibussowitsch 696d54fb17SJacob Faibussowitsch #endif // __cplusplus 706d54fb17SJacob Faibussowitsch 716d54fb17SJacob Faibussowitsch #endif // PETSC_DEVICE_CUPM_KERNELS_HPP 72