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 3490585354SJacob Faibussowitsch namespace functors 3590585354SJacob Faibussowitsch { 3690585354SJacob Faibussowitsch 3790585354SJacob Faibussowitsch template <typename T> 3890585354SJacob Faibussowitsch class plus_equals { 3990585354SJacob Faibussowitsch public: 4090585354SJacob Faibussowitsch using value_type = T; 4190585354SJacob Faibussowitsch 4290585354SJacob Faibussowitsch PETSC_HOSTDEVICE_DECL constexpr explicit plus_equals(value_type v = value_type{}) noexcept : v_{std::move(v)} { } 4390585354SJacob Faibussowitsch 4490585354SJacob Faibussowitsch PETSC_NODISCARD PETSC_HOSTDEVICE_INLINE_DECL constexpr value_type operator()(const value_type &val) const noexcept { return val + v_; } 4590585354SJacob Faibussowitsch 4690585354SJacob Faibussowitsch private: 4790585354SJacob Faibussowitsch value_type v_; 4890585354SJacob Faibussowitsch }; 4990585354SJacob Faibussowitsch 502ea277ceSJacob Faibussowitsch template <typename T> 512ea277ceSJacob Faibussowitsch class times_equals { 522ea277ceSJacob Faibussowitsch public: 532ea277ceSJacob Faibussowitsch using value_type = T; 542ea277ceSJacob Faibussowitsch 552ea277ceSJacob Faibussowitsch PETSC_HOSTDEVICE_DECL constexpr explicit times_equals(value_type v = value_type{}) noexcept : v_{std::move(v)} { } 562ea277ceSJacob Faibussowitsch 572ea277ceSJacob Faibussowitsch PETSC_NODISCARD PETSC_HOSTDEVICE_INLINE_DECL constexpr value_type operator()(const value_type &val) const noexcept { return val * v_; } 582ea277ceSJacob Faibussowitsch 592ea277ceSJacob Faibussowitsch private: 602ea277ceSJacob Faibussowitsch value_type v_; 612ea277ceSJacob Faibussowitsch }; 622ea277ceSJacob Faibussowitsch 63*025e0618SJacob Faibussowitsch template <typename T> 64*025e0618SJacob Faibussowitsch class axpy { 65*025e0618SJacob Faibussowitsch public: 66*025e0618SJacob Faibussowitsch using value_type = T; 67*025e0618SJacob Faibussowitsch 68*025e0618SJacob Faibussowitsch PETSC_HOSTDEVICE_DECL constexpr explicit axpy(value_type v = value_type{}) noexcept : v_{std::move(v)} { } 69*025e0618SJacob Faibussowitsch 70*025e0618SJacob Faibussowitsch PETSC_NODISCARD PETSC_HOSTDEVICE_INLINE_DECL constexpr value_type operator()(const value_type &x, const value_type &y) const noexcept { return v_ * x + y; } 71*025e0618SJacob Faibussowitsch 72*025e0618SJacob Faibussowitsch private: 73*025e0618SJacob Faibussowitsch value_type v_; 74*025e0618SJacob Faibussowitsch }; 75*025e0618SJacob Faibussowitsch 7690585354SJacob Faibussowitsch namespace 7790585354SJacob Faibussowitsch { 7890585354SJacob Faibussowitsch 7990585354SJacob Faibussowitsch template <typename T> 8090585354SJacob Faibussowitsch PETSC_HOSTDEVICE_INLINE_DECL constexpr plus_equals<T> make_plus_equals(const T &v) noexcept 8190585354SJacob Faibussowitsch { 8290585354SJacob Faibussowitsch return plus_equals<T>{v}; 8390585354SJacob Faibussowitsch } 8490585354SJacob Faibussowitsch 852ea277ceSJacob Faibussowitsch template <typename T> 862ea277ceSJacob Faibussowitsch PETSC_HOSTDEVICE_INLINE_DECL constexpr times_equals<T> make_times_equals(const T &v) noexcept 872ea277ceSJacob Faibussowitsch { 882ea277ceSJacob Faibussowitsch return times_equals<T>{v}; 892ea277ceSJacob Faibussowitsch } 902ea277ceSJacob Faibussowitsch 91*025e0618SJacob Faibussowitsch template <typename T> 92*025e0618SJacob Faibussowitsch PETSC_HOSTDEVICE_INLINE_DECL constexpr axpy<T> make_axpy(const T &v) noexcept 93*025e0618SJacob Faibussowitsch { 94*025e0618SJacob Faibussowitsch return axpy<T>{v}; 95*025e0618SJacob Faibussowitsch } 96*025e0618SJacob Faibussowitsch 9790585354SJacob Faibussowitsch } // anonymous namespace 9890585354SJacob Faibussowitsch 9990585354SJacob Faibussowitsch } // namespace functors 10090585354SJacob Faibussowitsch 1016d54fb17SJacob Faibussowitsch } // namespace cupm 1026d54fb17SJacob Faibussowitsch 1036d54fb17SJacob Faibussowitsch } // namespace device 1046d54fb17SJacob Faibussowitsch 1056d54fb17SJacob Faibussowitsch } // namespace Petsc 1066d54fb17SJacob Faibussowitsch 1076d54fb17SJacob Faibussowitsch #endif // __cplusplus 1086d54fb17SJacob Faibussowitsch 1096d54fb17SJacob Faibussowitsch #endif // PETSC_DEVICE_CUPM_KERNELS_HPP 110