From 573f2edd2cc2e44bbc6d797cca2924584bf01da4 Mon Sep 17 00:00:00 2001 From: Bruno Turcksin Date: Wed, 23 Nov 2022 13:20:13 -0500 Subject: [PATCH] Replace DEAL_II_CUDA_HOST_DEV with DEAL_II_HOST_DEV --- include/deal.II/base/numbers.h | 25 +- include/deal.II/base/point.h | 88 +++---- include/deal.II/base/tensor.h | 222 +++++++++--------- .../matrix_free/hanging_nodes_internal.h | 10 +- 4 files changed, 167 insertions(+), 178 deletions(-) diff --git a/include/deal.II/base/numbers.h b/include/deal.II/base/numbers.h index 930154200a..d53070eb1e 100644 --- a/include/deal.II/base/numbers.h +++ b/include/deal.II/base/numbers.h @@ -25,16 +25,15 @@ # include #endif +#include + #include #include #include #include -#ifdef DEAL_II_COMPILER_CUDA_AWARE -# define DEAL_II_CUDA_HOST_DEV __host__ __device__ -#else -# define DEAL_II_CUDA_HOST_DEV -#endif +#define DEAL_II_HOST_DEV KOKKOS_FUNCTION +#define DEAL_II_CUDA_HOST_DEV DEAL_II_HOST_DEV // Forward-declare the automatic differentiation types so we can add prototypes // for our own wrappers. @@ -269,7 +268,7 @@ namespace numbers /** * Check whether the given type can be used in CUDA device code. - * If not, DEAL_II_CUDA_HOST_DEV needs to be disabled for functions + * If not, DEAL_II_HOST_DEV needs to be disabled for functions * that use this type. */ template @@ -457,8 +456,8 @@ namespace numbers * * @note This function can also be used in CUDA device code. */ - static constexpr DEAL_II_CUDA_HOST_DEV const number & - conjugate(const number &x); + static constexpr DEAL_II_HOST_DEV const number & + conjugate(const number &x); /** * Return the square of the absolute value of the given number. Since the @@ -469,7 +468,7 @@ namespace numbers * for this function. */ template - static constexpr DEAL_II_CUDA_HOST_DEV + static constexpr DEAL_II_HOST_DEV std::enable_if_t::value && is_cuda_compatible::value, real_type> @@ -586,7 +585,7 @@ namespace numbers template - constexpr DEAL_II_CUDA_HOST_DEV const number & + constexpr DEAL_II_HOST_DEV const number & NumberTraits::conjugate(const number &x) { return x; @@ -596,7 +595,7 @@ namespace numbers template template - constexpr DEAL_II_CUDA_HOST_DEV + constexpr DEAL_II_HOST_DEV std::enable_if_t::value && is_cuda_compatible::value, typename NumberTraits::real_type> @@ -729,7 +728,7 @@ namespace internal template struct NumberType { - static constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV const T & + static constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV const T & value(const T &t) { return t; @@ -744,7 +743,7 @@ namespace internal // Type T is constructible from F. template - static constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV T + static constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV T value(const F &f, std::enable_if_t::type, typename std::decay::type>::value && diff --git a/include/deal.II/base/point.h b/include/deal.II/base/point.h index cf9cefa168..a58b844ebf 100644 --- a/include/deal.II/base/point.h +++ b/include/deal.II/base/point.h @@ -116,13 +116,13 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV + DEAL_II_HOST_DEV Point(); /** * Convert a tensor to a point. */ - explicit DEAL_II_CUDA_HOST_DEV + explicit DEAL_II_HOST_DEV Point(const Tensor<1, dim, Number> &); /** @@ -133,7 +133,7 @@ public: * * @note This function can also be used in CUDA device code. */ - explicit DEAL_II_CUDA_HOST_DEV + explicit DEAL_II_HOST_DEV Point(const Number x); /** @@ -145,7 +145,7 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV + DEAL_II_HOST_DEV Point(const Number x, const Number y); /** @@ -157,7 +157,7 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV + DEAL_II_HOST_DEV Point(const Number x, const Number y, const Number z); /** @@ -175,15 +175,15 @@ public: * * @note This function can also be used in CUDA device code. */ - static DEAL_II_CUDA_HOST_DEV Point - unit_vector(const unsigned int i); + static DEAL_II_HOST_DEV Point + unit_vector(const unsigned int i); /** * Read access to the indexth coordinate. * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV Number + DEAL_II_HOST_DEV Number operator()(const unsigned int index) const; /** @@ -191,8 +191,8 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV Number & - operator()(const unsigned int index); + DEAL_II_HOST_DEV Number & + operator()(const unsigned int index); /** * Assignment operator from Tensor<1, dim, Number> with different underlying @@ -213,8 +213,8 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV Point - operator+(const Tensor<1, dim, Number> &) const; + DEAL_II_HOST_DEV Point + operator+(const Tensor<1, dim, Number> &) const; /** * Subtract two points, i.e., obtain the vector that connects the two. As @@ -225,8 +225,8 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV Tensor<1, dim, Number> - operator-(const Point &) const; + DEAL_II_HOST_DEV Tensor<1, dim, Number> + operator-(const Point &) const; /** * Subtract a difference vector (represented by a Tensor@<1,dim@>) from the @@ -236,16 +236,16 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV Point - operator-(const Tensor<1, dim, Number> &) const; + DEAL_II_HOST_DEV Point + operator-(const Tensor<1, dim, Number> &) const; /** * The opposite vector. * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV Point - operator-() const; + DEAL_II_HOST_DEV Point + operator-() const; /** * @} @@ -264,7 +264,7 @@ public: * @relatesalso EnableIfScalar */ template - DEAL_II_CUDA_HOST_DEV Point< + DEAL_II_HOST_DEV Point< dim, typename ProductType::type>::type> @@ -276,7 +276,7 @@ public: * @note This function can also be used in CUDA device code. */ template - DEAL_II_CUDA_HOST_DEV Point< + DEAL_II_HOST_DEV Point< dim, typename ProductType::type>::type> @@ -287,7 +287,7 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV Number + DEAL_II_HOST_DEV Number operator*(const Tensor<1, dim, Number> &p) const; /** @@ -302,7 +302,7 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV typename numbers::NumberTraits::real_type + DEAL_II_HOST_DEV typename numbers::NumberTraits::real_type square() const; /** @@ -312,7 +312,7 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV typename numbers::NumberTraits::real_type + DEAL_II_HOST_DEV typename numbers::NumberTraits::real_type distance(const Point &p) const; /** @@ -321,7 +321,7 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV typename numbers::NumberTraits::real_type + DEAL_II_HOST_DEV typename numbers::NumberTraits::real_type distance_square(const Point &p) const; /** @@ -345,14 +345,14 @@ public: // At least clang-3.7 requires us to have a user-defined constructor // and we can't use 'Point::Point () = default' here. template -inline DEAL_II_CUDA_HOST_DEV +inline DEAL_II_HOST_DEV Point::Point() // NOLINT {} template -inline DEAL_II_CUDA_HOST_DEV +inline DEAL_II_HOST_DEV Point::Point(const Tensor<1, dim, Number> &t) : Tensor<1, dim, Number>(t) {} @@ -360,7 +360,7 @@ Point::Point(const Tensor<1, dim, Number> &t) template -inline DEAL_II_CUDA_HOST_DEV +inline DEAL_II_HOST_DEV Point::Point(const Number x) { # ifndef __CUDA_ARCH__ @@ -388,7 +388,7 @@ Point::Point(const Number x) template -inline DEAL_II_CUDA_HOST_DEV +inline DEAL_II_HOST_DEV Point::Point(const Number x, const Number y) { # ifndef __CUDA_ARCH__ @@ -411,7 +411,7 @@ Point::Point(const Number x, const Number y) template -inline DEAL_II_CUDA_HOST_DEV +inline DEAL_II_HOST_DEV Point::Point(const Number x, const Number y, const Number z) { # ifndef __CUDA_ARCH__ @@ -457,8 +457,8 @@ inline Point::Point( template -inline DEAL_II_CUDA_HOST_DEV Point - Point::unit_vector(unsigned int i) +inline DEAL_II_HOST_DEV Point + Point::unit_vector(unsigned int i) { Point p; p[i] = 1.; @@ -467,7 +467,7 @@ inline DEAL_II_CUDA_HOST_DEV Point template -inline DEAL_II_CUDA_HOST_DEV Number +inline DEAL_II_HOST_DEV Number Point::operator()(const unsigned int index) const { # ifndef __CUDA_ARCH__ @@ -479,7 +479,7 @@ Point::operator()(const unsigned int index) const template -inline DEAL_II_CUDA_HOST_DEV Number & +inline DEAL_II_HOST_DEV Number & Point::operator()(const unsigned int index) { # ifndef __CUDA_ARCH__ @@ -502,7 +502,7 @@ Point::operator=(const Tensor<1, dim, OtherNumber> &p) template -inline DEAL_II_CUDA_HOST_DEV Point +inline DEAL_II_HOST_DEV Point Point::operator+(const Tensor<1, dim, Number> &p) const { Point tmp = *this; @@ -513,7 +513,7 @@ Point::operator+(const Tensor<1, dim, Number> &p) const template -inline DEAL_II_CUDA_HOST_DEV Tensor<1, dim, Number> +inline DEAL_II_HOST_DEV Tensor<1, dim, Number> Point::operator-(const Point &p) const { return (Tensor<1, dim, Number>(*this) -= p); @@ -522,7 +522,7 @@ Point::operator-(const Point &p) const template -inline DEAL_II_CUDA_HOST_DEV Point +inline DEAL_II_HOST_DEV Point Point::operator-(const Tensor<1, dim, Number> &p) const { Point tmp = *this; @@ -533,7 +533,7 @@ Point::operator-(const Tensor<1, dim, Number> &p) const template -inline DEAL_II_CUDA_HOST_DEV Point +inline DEAL_II_HOST_DEV Point Point::operator-() const { Point result; @@ -546,7 +546,7 @@ Point::operator-() const template template -inline DEAL_II_CUDA_HOST_DEV +inline DEAL_II_HOST_DEV Point::type>::type> @@ -562,7 +562,7 @@ inline DEAL_II_CUDA_HOST_DEV template template -inline DEAL_II_CUDA_HOST_DEV +inline DEAL_II_HOST_DEV Point::type>::type> @@ -579,7 +579,7 @@ inline DEAL_II_CUDA_HOST_DEV template -inline DEAL_II_CUDA_HOST_DEV Number +inline DEAL_II_HOST_DEV Number Point::operator*(const Tensor<1, dim, Number> &p) const { Number res = Number(); @@ -590,7 +590,7 @@ Point::operator*(const Tensor<1, dim, Number> &p) const template -inline DEAL_II_CUDA_HOST_DEV typename numbers::NumberTraits::real_type +inline DEAL_II_HOST_DEV typename numbers::NumberTraits::real_type Point::square() const { return this->norm_square(); @@ -599,7 +599,7 @@ Point::square() const template -inline DEAL_II_CUDA_HOST_DEV typename numbers::NumberTraits::real_type +inline DEAL_II_HOST_DEV typename numbers::NumberTraits::real_type Point::distance(const Point &p) const { return std::sqrt(distance_square(p)); @@ -608,7 +608,7 @@ Point::distance(const Point &p) const template -inline DEAL_II_CUDA_HOST_DEV typename numbers::NumberTraits::real_type +inline DEAL_II_HOST_DEV typename numbers::NumberTraits::real_type Point::distance_square(const Point &p) const { Number sum = internal::NumberType::value(0.0); @@ -647,7 +647,7 @@ Point::serialize(Archive &ar, const unsigned int) * @relates Point */ template -inline DEAL_II_CUDA_HOST_DEV +inline DEAL_II_HOST_DEV Point::type>::type> diff --git a/include/deal.II/base/tensor.h b/include/deal.II/base/tensor.h index 7b21b36af4..2232675e72 100644 --- a/include/deal.II/base/tensor.h +++ b/include/deal.II/base/tensor.h @@ -139,7 +139,7 @@ public: * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV + constexpr DEAL_II_HOST_DEV Tensor(); /** @@ -150,7 +150,7 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV + constexpr DEAL_II_HOST_DEV Tensor(const Tensor<0, dim, OtherNumber> &initializer); /** @@ -159,20 +159,20 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV + constexpr DEAL_II_HOST_DEV Tensor(const OtherNumber &initializer); #ifdef DEAL_II_DELETED_MOVE_CONSTRUCTOR_BUG /** * Copy constructor */ - constexpr DEAL_II_CUDA_HOST_DEV + constexpr DEAL_II_HOST_DEV Tensor(const Tensor<0, dim, Number> &other); /** * Move constructor */ - constexpr DEAL_II_CUDA_HOST_DEV + constexpr DEAL_II_HOST_DEV Tensor(Tensor<0, dim, Number> &&other) noexcept; #endif @@ -234,7 +234,7 @@ public: * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV + constexpr DEAL_II_HOST_DEV operator Number &(); /** @@ -245,7 +245,7 @@ public: * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV operator const Number &() const; + constexpr DEAL_II_HOST_DEV operator const Number &() const; /** * Assignment from tensors with different underlying scalar type. This @@ -255,8 +255,8 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & - operator=(const Tensor<0, dim, OtherNumber> &rhs); + constexpr DEAL_II_HOST_DEV Tensor & + operator=(const Tensor<0, dim, OtherNumber> &rhs); #if defined(__INTEL_COMPILER) || defined(DEAL_II_DELETED_MOVE_CONSTRUCTOR_BUG) /** @@ -267,16 +267,16 @@ public: * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV Tensor & - operator=(const Tensor<0, dim, Number> &rhs); + constexpr DEAL_II_HOST_DEV Tensor & + operator=(const Tensor<0, dim, Number> &rhs); #endif #ifdef DEAL_II_DELETED_MOVE_CONSTRUCTOR_BUG /** * Move assignment operator */ - constexpr DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number> & - operator=(Tensor<0, dim, Number> &&other) noexcept; + constexpr DEAL_II_HOST_DEV Tensor<0, dim, Number> & + operator=(Tensor<0, dim, Number> &&other) noexcept; #endif /** @@ -286,8 +286,8 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & - operator=(const OtherNumber &d); + constexpr DEAL_II_HOST_DEV Tensor & + operator=(const OtherNumber &d); /** * Test for equality of two tensors. @@ -309,8 +309,8 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & - operator+=(const Tensor<0, dim, OtherNumber> &rhs); + constexpr DEAL_II_HOST_DEV Tensor & + operator+=(const Tensor<0, dim, OtherNumber> &rhs); /** * Subtract another scalar. @@ -318,8 +318,8 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & - operator-=(const Tensor<0, dim, OtherNumber> &rhs); + constexpr DEAL_II_HOST_DEV Tensor & + operator-=(const Tensor<0, dim, OtherNumber> &rhs); /** * Multiply the scalar with a factor. @@ -327,8 +327,8 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & - operator*=(const OtherNumber &factor); + constexpr DEAL_II_HOST_DEV Tensor & + operator*=(const OtherNumber &factor); /** * Divide the scalar by factor. @@ -336,15 +336,15 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & - operator/=(const OtherNumber &factor); + constexpr DEAL_II_HOST_DEV Tensor & + operator/=(const OtherNumber &factor); /** * Tensor with inverted entries. * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV Tensor + constexpr DEAL_II_HOST_DEV Tensor operator-() const; /** @@ -376,7 +376,7 @@ public: * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV real_type + constexpr DEAL_II_HOST_DEV real_type norm_square() const; /** @@ -547,7 +547,7 @@ public: * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV + constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor(); /** @@ -555,8 +555,7 @@ public: * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV explicit Tensor( - const array_type &initializer); + constexpr DEAL_II_HOST_DEV explicit Tensor(const array_type &initializer); /** * A constructor where the data is copied from an ArrayView object. @@ -572,7 +571,7 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV explicit Tensor( + constexpr DEAL_II_HOST_DEV explicit Tensor( const ArrayView &initializer); /** @@ -583,7 +582,7 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV + constexpr DEAL_II_HOST_DEV Tensor(const Tensor &initializer); /** @@ -617,16 +616,16 @@ public: * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV value_type & - operator[](const unsigned int i); + constexpr DEAL_II_HOST_DEV value_type & + operator[](const unsigned int i); /** * Read-only access operator. * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV const value_type & - operator[](const unsigned int i) const; + constexpr DEAL_II_HOST_DEV const value_type & + operator[](const unsigned int i) const; /** * Read access using TableIndices indices @@ -672,7 +671,7 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & + constexpr DEAL_II_HOST_DEV Tensor & operator=(const Tensor &rhs); /** @@ -718,7 +717,7 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & + constexpr DEAL_II_HOST_DEV Tensor & operator+=(const Tensor &); /** @@ -727,7 +726,7 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & + constexpr DEAL_II_HOST_DEV Tensor & operator-=(const Tensor &); /** @@ -737,8 +736,8 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & - operator*=(const OtherNumber &factor); + constexpr DEAL_II_HOST_DEV Tensor & + operator*=(const OtherNumber &factor); /** * Scale the vector by 1/factor. @@ -746,15 +745,15 @@ public: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV Tensor & - operator/=(const OtherNumber &factor); + constexpr DEAL_II_HOST_DEV Tensor & + operator/=(const OtherNumber &factor); /** * Unary minus operator. Negate all entries of a tensor. * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV Tensor + constexpr DEAL_II_HOST_DEV Tensor operator-() const; /** @@ -779,7 +778,7 @@ public: * * @note This function can also be used in CUDA device code. */ - DEAL_II_CUDA_HOST_DEV + DEAL_II_HOST_DEV typename numbers::NumberTraits::real_type norm() const; @@ -789,9 +788,8 @@ public: * * @note This function can also be used in CUDA device code. */ - constexpr DEAL_II_CUDA_HOST_DEV - typename numbers::NumberTraits::real_type - norm_square() const; + constexpr DEAL_II_HOST_DEV typename numbers::NumberTraits::real_type + norm_square() const; /** * Fill a vector with all tensor elements. @@ -880,7 +878,7 @@ private: * @note This function can also be used in CUDA device code. */ template - constexpr DEAL_II_CUDA_HOST_DEV + constexpr DEAL_II_HOST_DEV Tensor(const ArrayLike &initializer, std::index_sequence); // Allow an arbitrary Tensor to access the underlying values. @@ -957,7 +955,7 @@ namespace internal template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number>::Tensor() // Some auto-differentiable numbers need explicit // zero initialization such as adtl::adouble. @@ -968,7 +966,7 @@ Tensor<0, dim, Number>::Tensor() template template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number>::Tensor(const OtherNumber &initializer) : value(internal::NumberType::value(initializer)) {} @@ -977,7 +975,7 @@ Tensor<0, dim, Number>::Tensor(const OtherNumber &initializer) template template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number>::Tensor(const Tensor<0, dim, OtherNumber> &p) : Tensor{p.value} {} @@ -985,7 +983,7 @@ Tensor<0, dim, Number>::Tensor(const Tensor<0, dim, OtherNumber> &p) # ifdef DEAL_II_DELETED_MOVE_CONSTRUCTOR_BUG template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number>::Tensor(const Tensor<0, dim, Number> &other) : value{other.value} {} @@ -993,7 +991,7 @@ Tensor<0, dim, Number>::Tensor(const Tensor<0, dim, Number> &other) template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number>::Tensor(Tensor<0, dim, Number> &&other) noexcept : value{std::move(other.value)} {} @@ -1037,7 +1035,7 @@ Tensor<0, dim, Number>::end_raw() const template -constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number>::operator Number &() { // We cannot use Assert inside a CUDA kernel @@ -1051,7 +1049,7 @@ Tensor<0, dim, Number>::operator Number &() template constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number>::operator const Number &() const + DEAL_II_HOST_DEV Tensor<0, dim, Number>::operator const Number &() const { // We cannot use Assert inside a CUDA kernel # ifndef __CUDA_ARCH__ @@ -1065,9 +1063,8 @@ constexpr inline DEAL_II_ALWAYS_INLINE template template -constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number> & - Tensor<0, dim, Number>::operator=(const Tensor<0, dim, OtherNumber> &p) +constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number> & +Tensor<0, dim, Number>::operator=(const Tensor<0, dim, OtherNumber> &p) { value = internal::NumberType::value(p.value); return *this; @@ -1076,9 +1073,8 @@ constexpr inline DEAL_II_ALWAYS_INLINE # if defined(__INTEL_COMPILER) || defined(DEAL_II_DELETED_MOVE_CONSTRUCTOR_BUG) template -constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number> & - Tensor<0, dim, Number>::operator=(const Tensor<0, dim, Number> &p) +constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number> & +Tensor<0, dim, Number>::operator=(const Tensor<0, dim, Number> &p) { value = p.value; return *this; @@ -1087,7 +1083,7 @@ constexpr inline DEAL_II_ALWAYS_INLINE # ifdef DEAL_II_DELETED_MOVE_CONSTRUCTOR_BUG template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number> & +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number> & Tensor<0, dim, Number>::operator=(Tensor<0, dim, Number> &&other) noexcept { value = std::move(other.value); @@ -1099,9 +1095,8 @@ Tensor<0, dim, Number>::operator=(Tensor<0, dim, Number> &&other) noexcept template template -constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number> & - Tensor<0, dim, Number>::operator=(const OtherNumber &d) +constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number> & +Tensor<0, dim, Number>::operator=(const OtherNumber &d) { value = internal::NumberType::value(d); return *this; @@ -1136,9 +1131,8 @@ Tensor<0, dim, Number>::operator!=(const Tensor<0, dim, OtherNumber> &p) const template template -constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number> & - Tensor<0, dim, Number>::operator+=(const Tensor<0, dim, OtherNumber> &p) +constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number> & +Tensor<0, dim, Number>::operator+=(const Tensor<0, dim, OtherNumber> &p) { value += p.value; return *this; @@ -1147,9 +1141,8 @@ constexpr inline DEAL_II_ALWAYS_INLINE template template -constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number> & - Tensor<0, dim, Number>::operator-=(const Tensor<0, dim, OtherNumber> &p) +constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number> & +Tensor<0, dim, Number>::operator-=(const Tensor<0, dim, OtherNumber> &p) { value -= p.value; return *this; @@ -1162,7 +1155,7 @@ namespace internal namespace ComplexWorkaround { template - constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV void + constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV void multiply_assign_scalar(Number &val, const OtherNumber &s) { val *= s; @@ -1170,7 +1163,7 @@ namespace internal # ifdef __CUDA_ARCH__ template - constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV void + constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV void multiply_assign_scalar(std::complex &, const OtherNumber &) { printf("This function is not implemented for std::complex!\n"); @@ -1183,9 +1176,8 @@ namespace internal template template -constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number> & - Tensor<0, dim, Number>::operator*=(const OtherNumber &s) +constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number> & +Tensor<0, dim, Number>::operator*=(const OtherNumber &s) { internal::ComplexWorkaround::multiply_assign_scalar(value, s); return *this; @@ -1195,7 +1187,7 @@ constexpr inline DEAL_II_ALWAYS_INLINE template template -constexpr inline DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number> & +constexpr inline DEAL_II_HOST_DEV Tensor<0, dim, Number> & Tensor<0, dim, Number>::operator/=(const OtherNumber &s) { value /= s; @@ -1204,7 +1196,7 @@ Tensor<0, dim, Number>::operator/=(const OtherNumber &s) template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV Tensor<0, dim, Number> +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, Number> Tensor<0, dim, Number>::operator-() const { return -value; @@ -1222,7 +1214,7 @@ Tensor<0, dim, Number>::norm() const template -constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE +constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE typename Tensor<0, dim, Number>::real_type Tensor<0, dim, Number>::norm_square() const { @@ -1293,7 +1285,7 @@ constexpr unsigned int Tensor<0, dim, Number>::n_independent_components; template template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor::Tensor(const ArrayLike &initializer, std::index_sequence) : values{Tensor(initializer[indices])...} @@ -1305,7 +1297,7 @@ Tensor::Tensor(const ArrayLike &initializer, template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor::Tensor() // We would like to use =default, but this causes compile errors with some // MSVC versions and internal compiler errors with -O1 in gcc 5.4. @@ -1315,7 +1307,7 @@ Tensor::Tensor() template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor::Tensor(const array_type &initializer) : Tensor(initializer, std::make_index_sequence{}) {} @@ -1324,7 +1316,7 @@ Tensor::Tensor(const array_type &initializer) template template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor::Tensor( const ArrayView &initializer) { @@ -1341,7 +1333,7 @@ Tensor::Tensor( template template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor::Tensor( const Tensor &initializer) : Tensor(initializer, std::make_index_sequence{}) @@ -1393,11 +1385,10 @@ namespace internal namespace TensorSubscriptor { template - constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV ArrayElementType & - subscript(ArrayElementType * values, - const unsigned int i, - std::integral_constant) + constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV ArrayElementType & + subscript(ArrayElementType * values, + const unsigned int i, + std::integral_constant) { // We cannot use Assert in a CUDA kernel # ifndef __CUDA_ARCH__ @@ -1407,11 +1398,10 @@ namespace internal } template - constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV ArrayElementType & - subscript(ArrayElementType *dummy, - const unsigned int, - std::integral_constant) + constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV ArrayElementType & + subscript(ArrayElementType *dummy, + const unsigned int, + std::integral_constant) { // We cannot use Assert in a CUDA kernel # ifndef __CUDA_ARCH__ @@ -1427,7 +1417,7 @@ namespace internal template -constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV typename Tensor::value_type & Tensor::operator[](const unsigned int i) { @@ -1438,7 +1428,7 @@ constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV template constexpr DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV const typename Tensor::value_type & + DEAL_II_HOST_DEV const typename Tensor::value_type & Tensor::operator[](const unsigned int i) const { # ifndef DEAL_II_COMPILER_CUDA_AWARE @@ -1606,7 +1596,7 @@ Tensor::operator!=( template template constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor & + DEAL_II_HOST_DEV Tensor & Tensor::operator+=( const Tensor &p) { @@ -1619,7 +1609,7 @@ constexpr inline DEAL_II_ALWAYS_INLINE template template constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor & + DEAL_II_HOST_DEV Tensor & Tensor::operator-=( const Tensor &p) { @@ -1632,7 +1622,7 @@ constexpr inline DEAL_II_ALWAYS_INLINE template template constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor & + DEAL_II_HOST_DEV Tensor & Tensor::operator*=(const OtherNumber &s) { for (unsigned int i = 0; i < dim; ++i) @@ -1654,7 +1644,7 @@ namespace internal typename ProductType::type>::value && !std::is_same::value, int> = 0> - constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE void + constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE void division_operator(Tensor (&t)[dim], const OtherNumber &factor) { @@ -1674,7 +1664,7 @@ namespace internal typename ProductType::type>::value || std::is_same::value, int> = 0> - constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE void + constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE void division_operator(Tensor (&t)[dim], const OtherNumber &factor) { @@ -1689,7 +1679,7 @@ namespace internal template template constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor & + DEAL_II_HOST_DEV Tensor & Tensor::operator/=(const OtherNumber &s) { internal::TensorImplementation::division_operator(values, s); @@ -1699,7 +1689,7 @@ constexpr inline DEAL_II_ALWAYS_INLINE template constexpr inline DEAL_II_ALWAYS_INLINE - DEAL_II_CUDA_HOST_DEV Tensor + DEAL_II_HOST_DEV Tensor Tensor::operator-() const { Tensor tmp; @@ -1722,7 +1712,7 @@ Tensor::norm() const template -constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr inline DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV typename numbers::NumberTraits::real_type Tensor::norm_square() const { @@ -1938,7 +1928,7 @@ operator<<(std::ostream &out, const Tensor<0, dim, Number> &p) * @relatesalso Tensor */ template -constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE +constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE typename ProductType::type operator*(const Other &object, const Tensor<0, dim, Number> &t) { @@ -1958,7 +1948,7 @@ constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE * @relatesalso Tensor */ template -constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE +constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE typename ProductType::type operator*(const Tensor<0, dim, Number> &t, const Other &object) { @@ -1978,7 +1968,7 @@ constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE * @relatesalso Tensor */ template -DEAL_II_CUDA_HOST_DEV constexpr DEAL_II_ALWAYS_INLINE +DEAL_II_HOST_DEV constexpr DEAL_II_ALWAYS_INLINE typename ProductType::type operator*(const Tensor<0, dim, Number> & src1, const Tensor<0, dim, OtherNumber> &src2) @@ -1996,7 +1986,7 @@ DEAL_II_CUDA_HOST_DEV constexpr DEAL_II_ALWAYS_INLINE * @relatesalso Tensor */ template -DEAL_II_CUDA_HOST_DEV constexpr DEAL_II_ALWAYS_INLINE +DEAL_II_HOST_DEV constexpr DEAL_II_ALWAYS_INLINE Tensor<0, dim, typename ProductType -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, typename ProductType::type> operator+(const Tensor<0, dim, Number> & p, const Tensor<0, dim, OtherNumber> &q) @@ -2032,7 +2022,7 @@ constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV * @relatesalso Tensor */ template -constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV +constexpr DEAL_II_ALWAYS_INLINE DEAL_II_HOST_DEV Tensor<0, dim, typename ProductType::type> operator-(const Tensor<0, dim, Number> & p, const Tensor<0, dim, OtherNumber> &q) @@ -2054,7 +2044,7 @@ constexpr DEAL_II_ALWAYS_INLINE DEAL_II_CUDA_HOST_DEV * @relatesalso Tensor */ template -constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE +constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE Tensor -DEAL_II_CUDA_HOST_DEV constexpr inline DEAL_II_ALWAYS_INLINE +DEAL_II_HOST_DEV constexpr inline DEAL_II_ALWAYS_INLINE Tensor::type, @@ -2106,7 +2096,7 @@ namespace internal !std::is_integral< typename ProductType::type>::value, int> = 0> - constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE + constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE Tensor::type> division_operator(const Tensor &t, const OtherNumber & factor) @@ -2128,7 +2118,7 @@ namespace internal std::is_integral< typename ProductType::type>::value, int> = 0> - constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE + constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE Tensor::type> division_operator(const Tensor &t, const OtherNumber & factor) @@ -2153,7 +2143,7 @@ namespace internal * @relatesalso Tensor */ template -constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE +constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE Tensor -constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE +constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE Tensor::type> operator+(const Tensor & p, const Tensor &q) @@ -2198,7 +2188,7 @@ constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE * @relatesalso Tensor */ template -constexpr DEAL_II_CUDA_HOST_DEV inline DEAL_II_ALWAYS_INLINE +constexpr DEAL_II_HOST_DEV inline DEAL_II_ALWAYS_INLINE Tensor::type> operator-(const Tensor & p, const Tensor &q) diff --git a/include/deal.II/matrix_free/hanging_nodes_internal.h b/include/deal.II/matrix_free/hanging_nodes_internal.h index 3942abed73..438eaba272 100644 --- a/include/deal.II/matrix_free/hanging_nodes_internal.h +++ b/include/deal.II/matrix_free/hanging_nodes_internal.h @@ -199,7 +199,7 @@ namespace internal * would be an integer which would in turn trigger a compiler warning when * we tried to assign it to an object of type UpdateFlags. */ - DEAL_II_CUDA_HOST_DEV inline ConstraintKinds + DEAL_II_HOST_DEV inline ConstraintKinds operator|(const ConstraintKinds f1, const ConstraintKinds f2) { return static_cast(static_cast(f1) | @@ -212,7 +212,7 @@ namespace internal * Global operator which sets the bits from the second argument also in the * first one. */ - DEAL_II_CUDA_HOST_DEV inline ConstraintKinds & + DEAL_II_HOST_DEV inline ConstraintKinds & operator|=(ConstraintKinds &f1, const ConstraintKinds f2) { f1 = f1 | f2; @@ -224,7 +224,7 @@ namespace internal /** * Global operator which checks inequality. */ - DEAL_II_CUDA_HOST_DEV inline bool + DEAL_II_HOST_DEV inline bool operator!=(const ConstraintKinds f1, const ConstraintKinds f2) { return static_cast(f1) != static_cast(f2); @@ -236,7 +236,7 @@ namespace internal * Global operator which checks if the first argument is less than the * second. */ - DEAL_II_CUDA_HOST_DEV inline bool + DEAL_II_HOST_DEV inline bool operator<(const ConstraintKinds f1, const ConstraintKinds f2) { return static_cast(f1) < static_cast(f2); @@ -247,7 +247,7 @@ namespace internal /** * Global operator which performs a binary and for the provided arguments. */ - DEAL_II_CUDA_HOST_DEV inline ConstraintKinds + DEAL_II_HOST_DEV inline ConstraintKinds operator&(const ConstraintKinds f1, const ConstraintKinds f2) { return static_cast(static_cast(f1) & -- 2.39.5