diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h
index 4dae494a8a..9cc131ee62 100644
--- a/dace/runtime/include/dace/math.h
+++ b/dace/runtime/include/dace/math.h
@@ -496,7 +496,7 @@ namespace dace
 {
     namespace math
     {       
-        static DACE_CONSTEXPR typeless_pi pi{};
+        static DACE_CONSTEXPR DACE_HostDev typeless_pi pi{};
         static DACE_CONSTEXPR typeless_nan nan{};
         //////////////////////////////////////////////////////
         template<typename T>
diff --git a/dace/runtime/include/dace/nan.h b/dace/runtime/include/dace/nan.h
index a8d1eb4c52..b4bac93980 100644
--- a/dace/runtime/include/dace/nan.h
+++ b/dace/runtime/include/dace/nan.h
@@ -13,101 +13,119 @@ namespace dace
         // Defines a typeless Pi
         struct typeless_nan
         {
+            DACE_CONSTEXPR DACE_HDFI typeless_nan() noexcept = default;
+            DACE_CONSTEXPR DACE_HDFI typeless_nan(const typeless_nan&) noexcept = default;
+            DACE_CONSTEXPR DACE_HDFI typeless_nan(typeless_nan&&) noexcept = default;
+            DACE_HDFI ~typeless_nan() noexcept = default;
+
+#ifndef DACE_XILINX
+            DACE_CONSTEXPR DACE_HDFI typeless_nan& operator=(const typeless_nan&) noexcept = default;
+            DACE_CONSTEXPR DACE_HDFI typeless_nan& operator=(typeless_nan&&) noexcept = default;
+#endif
+
             operator int() const = delete;
-            operator float() const
+            DACE_CONSTEXPR DACE_HDFI operator float() const
             {
                 return std::numeric_limits<float>::quiet_NaN();
             }
-            operator double() const
+            DACE_CONSTEXPR DACE_HDFI operator double() const
             {
                 return std::numeric_limits<double>::quiet_NaN();
             }
-            operator long double() const
+
+#if !( defined(__CUDACC__) || defined(__HIPCC__) )
+            //There is no long double on the GPU
+            DACE_CONSTEXPR DACE_HDFI operator long double() const
             {
                 return std::numeric_limits<long double>::quiet_NaN();
             }
-            typeless_nan operator+() const
+#endif
+            DACE_CONSTEXPR DACE_HDFI typeless_nan operator+() const
             {
                 return typeless_nan{};
             }
-            typeless_nan operator-() const
+            DACE_CONSTEXPR DACE_HDFI typeless_nan operator-() const
             {
                 return typeless_nan{};
             }
         };
 
         template<typename T>
-        DACE_CONSTEXPR typename std::enable_if<std::is_floating_point<T>::value, typeless_nan>::type
-        operator*(const T&,  const typeless_nan&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI std::enable_if_t<std::is_floating_point<T>::value || std::is_integral<T>::value, typeless_nan>
+        operator*(const T&,  const typeless_nan&) noexcept { return typeless_nan{}; }
 
         template<typename T>
-        DACE_CONSTEXPR typename std::enable_if<std::is_floating_point<T>::value, typeless_nan>::type
-        operator*(const typeless_nan&,  const T&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI std::enable_if_t<std::is_floating_point<T>::value || std::is_integral<T>::value, typeless_nan>
+        operator*(const typeless_nan&,  const T&) noexcept { return typeless_nan{}; }
 
-        inline typeless_nan
-        operator*(const typeless_nan&,  const typeless_nan&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI typeless_nan
+        operator*(const typeless_nan&,  const typeless_nan&) noexcept { return typeless_nan{}; }
 
 
         template<typename T>
-        DACE_CONSTEXPR typename std::enable_if<std::is_floating_point<T>::value, typeless_nan>::type
-        operator+(const T&,  const typeless_nan&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI std::enable_if_t<std::is_floating_point<T>::value || std::is_integral<T>::value, typeless_nan>
+        operator+(const T&,  const typeless_nan&) noexcept { return typeless_nan{}; }
 
         template<typename T>
-        DACE_CONSTEXPR typename std::enable_if<std::is_floating_point<T>::value, typeless_nan>::type
-        operator+(const typeless_nan&,  const T&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI std::enable_if_t<std::is_floating_point<T>::value || std::is_integral<T>::value, typeless_nan>
+        operator+(const typeless_nan&,  const T&) noexcept { return typeless_nan{}; }
 
-        inline typeless_nan
-        operator+(const typeless_nan&,  const typeless_nan&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI typeless_nan
+        operator+(const typeless_nan&,  const typeless_nan&) noexcept { return typeless_nan{}; }
 
 
         template<typename T>
-        DACE_CONSTEXPR typename std::enable_if<std::is_floating_point<T>::value, typeless_nan>::type
-        operator-(const T&,  const typeless_nan&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI std::enable_if_t<std::is_floating_point<T>::value || std::is_integral<T>::value, typeless_nan>
+        operator-(const T&,  const typeless_nan&) noexcept { return typeless_nan{}; }
 
         template<typename T>
-        DACE_CONSTEXPR typename std::enable_if<std::is_floating_point<T>::value, typeless_nan>::type
-        operator-(const typeless_nan&,  const T&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI std::enable_if_t<std::is_floating_point<T>::value || std::is_integral<T>::value, typeless_nan>
+        operator-(const typeless_nan&,  const T&) noexcept { return typeless_nan{}; }
 
-        inline typeless_nan
-        operator-(const typeless_nan&,  const typeless_nan&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI typeless_nan
+        operator-(const typeless_nan&,  const typeless_nan&) noexcept { return typeless_nan{}; }
 
 
         template<typename T>
-        DACE_CONSTEXPR typename std::enable_if<std::is_floating_point<T>::value, typeless_nan>::type
-        operator/(const T&,  const typeless_nan&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI std::enable_if_t<std::is_floating_point<T>::value || std::is_integral<T>::value, typeless_nan>
+        operator/(const T&,  const typeless_nan&) noexcept { return typeless_nan{}; }
 
         template<typename T>
-        DACE_CONSTEXPR typename std::enable_if<std::is_floating_point<T>::value, typeless_nan>::type
-        operator/(const typeless_nan&,  const T&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI std::enable_if_t<std::is_floating_point<T>::value || std::is_integral<T>::value, typeless_nan>
+        operator/(const typeless_nan&,  const T&) noexcept { return typeless_nan{}; }
 
-        inline typeless_nan
-        operator/(const typeless_nan&,  const typeless_nan&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI typeless_nan
+        operator/(const typeless_nan&,  const typeless_nan&) noexcept { return typeless_nan{}; }
 
 
         template<typename T>
-        DACE_CONSTEXPR typename std::enable_if<std::is_floating_point<T>::value, typeless_nan>::type
-        operator%(const T&,  const typeless_nan&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI std::enable_if_t<std::is_floating_point<T>::value || std::is_integral<T>::value, typeless_nan>
+        operator%(const T&,  const typeless_nan&) noexcept { return typeless_nan{}; }
 
         template<typename T>
-        DACE_CONSTEXPR typename std::enable_if<std::is_floating_point<T>::value, typeless_nan>::type
-        operator%(const typeless_nan&,  const T&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI std::enable_if_t<std::is_floating_point<T>::value || std::is_integral<T>::value, typeless_nan>
+        operator%(const typeless_nan&,  const T&) noexcept { return typeless_nan{}; }
 
-        inline typeless_nan
-        operator%(const typeless_nan&,  const typeless_nan&) { return typeless_nan{}; }
+        DACE_CONSTEXPR DACE_HDFI typeless_nan
+        operator%(const typeless_nan&,  const typeless_nan&) noexcept { return typeless_nan{}; }
 
-    }
-}
+        DACE_HDFI typeless_nan ipow(const typeless_nan&, const unsigned int&) {
+            return typeless_nan{};
+        }
 
 	//These functions allows to perfrom operations with `typeless_nan` instances.
-#	define FADAPT(F) DACE_CONSTEXPR ::dace::math::typeless_nan F (::dace::math::typeless_nan) { return ::dace::math::typeless_nan{}; }
-#	define FADAPT2(F) template<typename T1> DACE_CONSTEXPR dace::math::typeless_nan F (T1&&, dace::math::typeless_nan) { return ::dace::math::typeless_nan{}; }; \
-			  template<typename T2> DACE_CONSTEXPR dace::math::typeless_nan F (dace::math::typeless_nan, T2&&) { return ::dace::math::typeless_nan{}; }; \
-			  DACE_CONSTEXPR ::dace::math::typeless_nan F (dace::math::typeless_nan, dace::math::typeless_nan) { return ::dace::math::typeless_nan{}; }
+#	define FADAPT(F) DACE_CONSTEXPR DACE_HDFI typeless_nan F (const typeless_nan&) noexcept { return typeless_nan{}; }
+#	define FADAPT2(F) template<typename T1> DACE_CONSTEXPR DACE_HDFI typeless_nan F (T1&&, dace::math::typeless_nan) noexcept { return typeless_nan{}; }; \
+			  template<typename T2> DACE_CONSTEXPR DACE_HDFI typeless_nan F (const typeless_nan&, T2&&) noexcept { return typeless_nan{}; }; \
+			  DACE_CONSTEXPR DACE_HDFI typeless_nan F (const typeless_nan&, const typeless_nan&) noexcept { return typeless_nan{}; }
         FADAPT(tanh); FADAPT(cos); FADAPT(sin); FADAPT(sqrt); FADAPT(tan);
         FADAPT(acos); FADAPT(asin); FADAPT(atan); FADAPT(log); FADAPT(exp);
         FADAPT(floor); FADAPT(ceil); FADAPT(round); FADAPT(abs);
         FADAPT2(max); FADAPT2(min);
 #       undef FADAPT2
 #	undef FADAPT
+    }
+}
+
 
 #endif  // __DACE_NAN_H
diff --git a/dace/runtime/include/dace/pi.h b/dace/runtime/include/dace/pi.h
index 331b8c1636..818a22f6d0 100644
--- a/dace/runtime/include/dace/pi.h
+++ b/dace/runtime/include/dace/pi.h
@@ -2,6 +2,8 @@
 #ifndef __DACE_PI_H
 #define __DACE_PI_H
 
+#include <type_traits>
+
 // Classes that are used to define a typeless Pi
 
 //#define _USE_MATH_DEFINES
@@ -16,233 +18,274 @@ namespace dace
     {
         //////////////////////////////////////////////////////
         // Defines a typeless Pi
-        struct typeless_pi
+
+        template<typename T>
+        struct is_typeless_pi { static constexpr bool value = false; };
+        #define MAKE_TYPELESS_PI(type) template<> struct is_typeless_pi<type> { static constexpr bool value = true; }
+
+        struct typeless_pi;
+
+        /* Represents $m * \pi$. */
+        struct typeless_pi_mult
         {
-            double value() const { return M_PI; }
-            operator int() const
-            {
-                return int(this->value());
-            }
-            operator float() const
-            {
-                return float(this->value());
-            }
-            operator double() const
-            {
-                return double(this->value());
-            }
+            int mult;
+
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_mult(int m): mult(m) {}
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_mult() noexcept: typeless_pi_mult(1) {};
+
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_mult(const typeless_pi&) noexcept: typeless_pi_mult(1) {};
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_mult(const typeless_pi_mult&) noexcept = default;
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_mult(typeless_pi_mult&&) noexcept = default;
+            DACE_HDFI ~typeless_pi_mult() noexcept = default;
+
+#ifndef DACE_XILINX
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_mult& operator=(const typeless_pi_mult&) noexcept = default;
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_mult& operator=(typeless_pi_mult&&) noexcept = default;
+#endif
+
+            template<
+                typename T,
+                typename = std::enable_if_t<std::is_integral<T>::value>
+            >
+            DACE_CONSTEXPR DACE_HDFI operator T() const noexcept
+            { return T(mult * M_PI); }
+
+            DACE_CONSTEXPR DACE_HDFI operator float() const noexcept
+            { return float(mult * M_PI); }
+
+            DACE_CONSTEXPR DACE_HDFI operator double() const noexcept
+            { return mult * M_PI; }
+
+            DACE_CONSTEXPR DACE_HDFI operator long double() const noexcept
+            { return (long double)(mult * M_PI); }
+
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator+() const noexcept
+            { return *this; }
+
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator-() const noexcept
+            { return typeless_pi_mult(-this->mult); }
         };
-        struct typeless_pi_mult : typeless_pi
+        MAKE_TYPELESS_PI(typeless_pi_mult);
+
+        /* Represents $\pi$ */
+        struct typeless_pi
         {
-            int mult; typeless_pi_mult(int m = 1) : mult(m) {}
-            double value() const { return mult * M_PI; }
-
-            operator int() const
-            {
-                return int(this->value());
-            }
-            operator float() const
-            {
-                return float(this->value());
-            }
-            operator double() const
-            {
-                return double(this->value());
-            }
+            DACE_CONSTEXPR DACE_HDFI typeless_pi() noexcept = default;
+            DACE_CONSTEXPR DACE_HDFI typeless_pi(const typeless_pi&) noexcept = default;
+            DACE_CONSTEXPR DACE_HDFI typeless_pi(typeless_pi&&) noexcept = default;
+            DACE_HDFI ~typeless_pi() noexcept = default;
+#ifndef DACE_XILINX
+            DACE_CONSTEXPR DACE_HDFI typeless_pi& operator=(const typeless_pi&) noexcept = default;
+            DACE_CONSTEXPR DACE_HDFI typeless_pi& operator=(typeless_pi&&) noexcept = default;
+#endif
+
+            template<
+                typename T,
+                typename = std::enable_if_t<std::is_integral<T>::value>
+            >
+            DACE_CONSTEXPR DACE_HDFI operator T() const noexcept
+            { return T(M_PI); }
+
+            DACE_CONSTEXPR DACE_HDFI operator float() const noexcept
+            { return float(M_PI); }
+
+            DACE_CONSTEXPR DACE_HDFI operator double() const noexcept
+            { return M_PI; }
+
+            DACE_CONSTEXPR DACE_HDFI operator long double() const noexcept
+            { return (long double)(M_PI); }
+
+            DACE_CONSTEXPR DACE_HDFI typeless_pi operator+() const noexcept
+            { return *this; }
+
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator-() const noexcept
+            { return typeless_pi_mult(-1); }
         };
-        struct typeless_pi_exp : typeless_pi_mult
+        MAKE_TYPELESS_PI(typeless_pi);
+
+        /* Represents $m * \pi^{e}$ */
+        struct typeless_pi_exp
         {
-            int mult, exp; typeless_pi_exp(int m = 1, int e = 1) : mult(m), exp(e) {}
-            double value() const { return mult * std::pow(M_PI, exp); }
-            operator int() const
-            {
-                return int(this->value());
-            }
-            operator float() const
-            {
-                return float(this->value());
-            }
-            operator double() const
-            {
-                return double(this->value());
-            }
+            int mult, exp;
+
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_exp(int m, int e): mult(m), exp(e) {}
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_exp() noexcept: typeless_pi_exp(1, 1) {};
+
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_exp(const typeless_pi_exp&) noexcept = default;
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_exp(typeless_pi_exp&&) noexcept = default;
+            DACE_HDFI ~typeless_pi_exp() noexcept = default;
+
+#ifndef DACE_XILINX
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_exp& operator=(const typeless_pi_exp&) noexcept = default;
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_exp& operator=(typeless_pi_exp&&) noexcept = default;
+#endif
+
+            template<
+                typename T,
+                typename = std::enable_if_t<std::is_integral<T>::value>
+            >
+            DACE_CONSTEXPR DACE_HDFI operator T() const noexcept
+            { return T(mult * std::pow(static_cast<T>(M_PI), exp)); }
+
+
+            /* We have to do the selection this way, because it seems as nvidia does
+             *  not provide `powl` and `powf` in the std namespace */
+            DACE_CONSTEXPR DACE_HDFI operator float() const
+            { using std::pow; return mult * pow(static_cast<float>(M_PI), exp); }
+
+            DACE_CONSTEXPR DACE_HDFI operator double() const
+            { using std::pow; return mult * std::pow(static_cast<double>(M_PI), exp); }
+
+#if !( defined(__CUDACC__) || defined(__HIPCC__) )
+            //There is no long double on the GPU
+            DACE_CONSTEXPR DACE_HDFI operator long double() const
+            { using std::pow; return mult * std::pow(static_cast<long double>(M_PI), exp); }
+#endif
+
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator+() const
+            { return *this; }
+
+            DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator-() const
+            { return typeless_pi_exp(-this->mult, this->exp); }
         };
-        inline typeless_pi_mult operator*(const typeless_pi&, const int& num)
-        {
-            return typeless_pi_mult(num);
-        }
-        inline typeless_pi_mult operator*(const typeless_pi_mult& p, const int& num)
-        {
-            return typeless_pi_mult(p.mult * num);
-        }
-        inline typeless_pi_exp operator*(const typeless_pi_exp& p, const int& num)
-        {
-            return typeless_pi_exp(p.mult * num, p.exp);
-        }
-        inline typeless_pi_mult operator*(const int& num, const typeless_pi&)
-        {
-            return typeless_pi_mult(num);
-        }
-        inline typeless_pi_mult operator*(const int& num, const typeless_pi_mult& p)
-        {
-            return typeless_pi_mult(num * p.mult);
-        }
-        inline typeless_pi_exp operator*(const int& num, const typeless_pi_exp& p)
-        {
-            return typeless_pi_exp(num * p.mult, p.exp);
-        }
-        template <typename T>
-        T operator+(const typeless_pi& p, const T& num)
-        {
-            return T(p.value()) + num;
-        }
-        template <typename T>
-        T operator-(const typeless_pi& p, const T& num)
-        {
-            return T(p.value()) - num;
-        }
+        MAKE_TYPELESS_PI(typeless_pi_exp);
 
-        template <typename T>
-        T operator*(const typeless_pi& p, const T& num)
-        {
-            return T(p.value()) * num;
-        }
-        template <typename T>
-        T operator/(const typeless_pi& p, const T& num)
-        {
-            return T(p.value()) / num;
-        }
-        template <typename T>
-        T operator+(const T& num, const typeless_pi& p)
-        {
-            return num + T(p.value());
-        }
-        template <typename T>
-        T operator-(const T& num, const typeless_pi& p)
-        {
-            return num - T(p.value());
-        }
-        template <typename T>
-        T operator*(const T& num, const typeless_pi& p)
-        {
-            return num * T(p.value());
-        }
-        template <typename T>
-        T operator/(const T& num, const typeless_pi& p)
-        {
-            return num / T(p.value());
-        }
-        template <typename T>
-        T operator+(const typeless_pi_mult& p, const T& num)
-        {
-            return T(p.value()) + num;
-        }
-        template <typename T>
-        T operator-(const typeless_pi_mult& p, const T& num)
-        {
-            return T(p.value()) - num;
-        }
 
-        template <typename T>
-        T operator*(const typeless_pi_mult& p, const T& num)
-        {
-            return T(p.value()) * num;
-        }
-        template <typename T>
-        T operator/(const typeless_pi_mult& p, const T& num)
-        {
-            return T(p.value()) / num;
-        }
-        template <typename T>
-        T operator+(const T& num, const typeless_pi_mult& p)
-        {
-            return num + T(p.value());
-        }
-        template <typename T>
-        T operator-(const T& num, const typeless_pi_mult& p)
-        {
-            return num - T(p.value());
-        }
-        template <typename T>
-        T operator*(const T& num, const typeless_pi_mult& p)
-        {
-            return num * T(p.value());
-        }
-        template <typename T>
-        T operator/(const T& num, const typeless_pi_mult& p)
-        {
-            return num / T(p.value());
-        }
-        template <typename T>
-        T operator+(const typeless_pi_exp& p, const T& num)
-        {
-            return T(p.value()) + num;
-        }
-        template <typename T>
-        T operator-(const typeless_pi_exp& p, const T& num)
-        {
-            return T(p.value()) - num;
-        }
+        DACE_CONSTEXPR DACE_HDFI int operator/(const typeless_pi&, const typeless_pi&) noexcept
+        { return 1; }
 
-        template <typename T>
-        T operator*(const typeless_pi_exp& p, const T& num)
-        {
-            return T(p.value()) * num;
-        }
-        template <typename T>
-        T operator/(const typeless_pi_exp& p, const T& num)
-        {
-            return T(p.value()) / num;
-        }
-        template <typename T>
-        T operator+(const T& num, const typeless_pi_exp& p)
-        {
-            return num + T(p.value());
-        }
-        template <typename T>
-        T operator-(const T& num, const typeless_pi_exp& p)
-        {
-            return num - T(p.value());
-        }
-        template <typename T>
-        T operator*(const T& num, const typeless_pi_exp& p)
-        {
-            return num * T(p.value());
-        }
-        template <typename T>
-        T operator/(const T& num, const typeless_pi_exp& p)
-        {
-            return num / T(p.value());
-        }
-        inline typeless_pi_mult operator-(const typeless_pi&)
-        {
-            return typeless_pi_mult(-1);
-        }
-        template <typename T>
-        typeless_pi_mult operator+(const typeless_pi&, const typeless_pi&)
-        {
-            return typeless_pi_mult(2);
-        }
-        template <typename T>
-        typeless_pi_mult operator+(const typeless_pi_mult& p1, const typeless_pi_mult& p2)
-        {
-            return typeless_pi_mult(p1.mult + p2.mult);
-        }
-        template <typename T>
-        typeless_pi_exp operator*(const typeless_pi_mult& p1, const typeless_pi_mult& p2)
-        {
-            return typeless_pi_exp(p1.mult * p2.mult, 2);
+        DACE_CONSTEXPR DACE_HDFI int operator-(const typeless_pi&, const typeless_pi&) noexcept
+        { return 0; }
+
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator*(const typeless_pi&, const int& num) noexcept
+        { return typeless_pi_mult(num); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator*(const int& num, const typeless_pi&) noexcept
+        { return typeless_pi_mult(num); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator*(const typeless_pi_mult& p, const int& num) noexcept
+        { return typeless_pi_mult(p.mult * num); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator*(const int& num, const typeless_pi_mult& p) noexcept
+        { return typeless_pi_mult(p.mult * num); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator+(const typeless_pi&, const typeless_pi&) noexcept
+        { return typeless_pi_mult(2); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator+(const typeless_pi&, const typeless_pi_mult& pi) noexcept
+        { return typeless_pi_mult(pi.mult + 1); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator+(const typeless_pi_mult& pi, const typeless_pi&) noexcept
+        { return typeless_pi_mult(pi.mult + 1); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator+(const typeless_pi_mult& pl, const typeless_pi_mult& pr) noexcept
+        { return typeless_pi_mult(pl.mult + pr.mult); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_mult operator-(const typeless_pi_mult& pl, const typeless_pi_mult& pr) noexcept
+        { return typeless_pi_mult(pl.mult - pr.mult); }
+
+        DACE_CONSTEXPR DACE_HDFI int operator/(const typeless_pi_mult& pl, const typeless_pi&) noexcept
+        { return pl.mult; }
+
+        DACE_CONSTEXPR DACE_HDFI double operator/(const typeless_pi& pl, const typeless_pi_mult& pr) noexcept
+        { return 1.0 / pr.mult; }
+
+
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator*(const typeless_pi&, const typeless_pi&) noexcept
+        { return typeless_pi_exp(1, 2); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator*(const typeless_pi_mult& pl, const typeless_pi_mult& pr) noexcept
+        { return typeless_pi_exp(pl.mult * pr.mult, 2); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator*(const typeless_pi_mult& pl, const typeless_pi&) noexcept
+        { return typeless_pi_exp(pl.mult, 2); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator*(const typeless_pi& pl, const typeless_pi_mult& pr) noexcept
+        { return typeless_pi_exp(pr.mult, 2); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator*(const typeless_pi_exp& pl, const typeless_pi_mult& pr) noexcept
+        { return typeless_pi_exp(pl.mult * pr.mult, pl.exp + 1); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator*(const typeless_pi_mult& pl, const typeless_pi_exp& pr) noexcept
+        { return pr * pl; }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator*(const typeless_pi_exp& pl, const typeless_pi_exp& pr) noexcept
+        { return typeless_pi_exp(pl.mult * pr.mult, pr.exp + pl.exp); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator*(const typeless_pi_exp& pl, const int& num) noexcept
+        { return typeless_pi_exp(pl.mult * num, pl.exp); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator*(const int& num, const typeless_pi_exp& pr) noexcept
+        { return typeless_pi_exp(pr.mult * num, pr.exp); }
+
+        DACE_CONSTEXPR DACE_HDFI typeless_pi_exp operator/(const typeless_pi_exp& pl, const typeless_pi&) noexcept
+        { return typeless_pi_exp(pl.mult, pl.exp - 1); }
+
+
+        // The code generator guarantees us that `b > 0`.
+        DACE_HDFI typeless_pi_exp ipow(const typeless_pi_mult& pi, const unsigned int& b) {
+            return typeless_pi_exp(pow(pi.mult, b), b);
         }
-        template <typename T>
-        typeless_pi_exp operator*(const typeless_pi&, const typeless_pi&)
-        {
-            return typeless_pi_exp(1, 2);
+        DACE_HDFI typeless_pi_exp ipow(const typeless_pi& pi, const unsigned int& b) {
+            return typeless_pi_exp(1, b);
         }
-        template <typename T>
-        typeless_pi_exp operator*(const typeless_pi_exp& p1, const typeless_pi_exp& p2)
-        {
-            return typeless_pi_exp(p1.mult * p2.mult, p1.exp + p2.exp);
+        DACE_HDFI typeless_pi_exp ipow(const typeless_pi_exp& pi, const unsigned int& b) {
+            return typeless_pi_exp(pow(pi.mult, b), pi.exp * b);
         }
+
+#       define DEF_PI_OPS(op) 										\
+	template<typename T, typename PI, typename = std::enable_if_t<is_typeless_pi<PI>::value && (!is_typeless_pi<T>::value)> >	\
+	DACE_CONSTEXPR DACE_HDFI T operator op (const T& lhs, const PI& pi) noexcept			\
+	{ return lhs op (static_cast<T>(pi)); }								\
+	template<typename PI, typename T, typename = std::enable_if_t<is_typeless_pi<PI>::value && (!is_typeless_pi<T>::value)> >	\
+	DACE_CONSTEXPR DACE_HDFI T operator op (const PI& pi, const T& rhs) noexcept			\
+	{ return (static_cast<T>(pi)) op rhs; }
+
+	DEF_PI_OPS(+);
+	DEF_PI_OPS(-);
+	DEF_PI_OPS(/);
+	DEF_PI_OPS(*);
+
+        DACE_CONSTEXPR DACE_HDFI int sin(const typeless_pi&) noexcept
+        { return 0; }
+
+        DACE_CONSTEXPR DACE_HDFI int sin(const typeless_pi_mult& pi) noexcept
+        { return 0; }
+
+	DACE_HDFI double sin(const typeless_pi_exp& pi) noexcept
+	{ return std::sin(static_cast<double>(pi)); }
+
+        DACE_CONSTEXPR DACE_HDFI int cos(const typeless_pi&) noexcept
+        { return 1; }
+
+        DACE_CONSTEXPR DACE_HDFI int cos(const typeless_pi_mult& pi) noexcept
+        { return (pi.mult % 2 == 0) ? 1 : (-1); }
+
+	DACE_HDFI double cos(const typeless_pi_exp& pi) noexcept
+	{ return std::cos(static_cast<double>(pi)); }
+
+
+#       define DEF_PI_TRIGO(F)						\
+	DACE_HDFI double F (const typeless_pi& pi) noexcept		\
+	{ return std:: F( static_cast<double>(pi) ); }			\
+	DACE_HDFI double F (const typeless_pi_mult& pi) noexcept	\
+	{ return std:: F( static_cast<double>(pi) ); }			\
+	DACE_HDFI double F (const typeless_pi_exp& pi) noexcept		\
+	{ return std:: F( static_cast<double>(pi) ); }
+
+        DEF_PI_TRIGO(asin);
+        DEF_PI_TRIGO(acos);
+        DEF_PI_TRIGO(tan);
+        DEF_PI_TRIGO(atan);
+        DEF_PI_TRIGO(exp);
+        DEF_PI_TRIGO(log);
+
+
+#       undef DEF_PI_TRIGO
+#       undef DEF_PI_OPS
+#	undef MAKE_TYPELESS_PI
     }
 }
 
diff --git a/dace/runtime/include/dace/types.h b/dace/runtime/include/dace/types.h
index e5eed1e35e..9a8676e0d4 100644
--- a/dace/runtime/include/dace/types.h
+++ b/dace/runtime/include/dace/types.h
@@ -49,7 +49,13 @@
     #define DACE_HDFI __host__ __device__ __forceinline__
     #define DACE_HFI __host__ __forceinline__
     #define DACE_DFI __device__ __forceinline__
+    #define DACE_HostDev __host__ __device__
+    #define DACE_Host __host__ 
+    #define DACE_Dev __device__
 #else
+    #define DACE_HostDev
+    #define DACE_Host 
+    #define DACE_Dev
     #define DACE_HDFI inline
     #define DACE_HFI inline
     #define DACE_DFI inline
diff --git a/tests/numpy/constants_test.py b/tests/numpy/constants_test.py
new file mode 100644
index 0000000000..6aa3b15621
--- /dev/null
+++ b/tests/numpy/constants_test.py
@@ -0,0 +1,170 @@
+# Copyright 2019-2024 ETH Zurich and the DaCe authors. All rights reserved.
+import dace
+import numpy as np
+import uuid
+import math
+import pytest
+
+def _make_sdfg(
+        code: str,
+        dtype = dace.float64,
+) -> dace.SDFG:
+    """Generates an SDFG that writes an expression to an array.
+    """
+    sdfg = dace.SDFG(name=f"const_test_{str(uuid.uuid1()).replace('-', '_')}")
+    state = sdfg.add_state(is_start_block=True)
+    sdfg.add_array(
+            "out",
+            shape=(10,),
+            dtype=dtype,
+            transient=False,
+    )
+
+    state.add_mapped_tasklet(
+            "comput",
+            map_ranges={"__i": "0:10"},
+            inputs={},
+            code=f"__out = {code}",
+            outputs={"__out": dace.Memlet("out[__i]")},
+            external_edges=True,
+    )
+    sdfg.validate()
+    return sdfg
+
+
+def _test_sdfg(
+        sdfg: dace.SDFG,
+        expected,
+        dtype = np.float64,
+):
+    out = np.zeros(10, dtype=dtype)
+    sdfg.apply_gpu_transformations()
+    sdfg(out=out)
+    assert np.allclose(out, expected, equal_nan=True), f"Expected {expected}, but got {out[0]}"
+
+
+def _perform_test(
+        code,
+        expected,
+        dtype = np.float64,
+):
+    print(f"PERFORM: {code}")
+    dace_dtype = dace.dtypes.dtype_to_typeclass(dtype)
+    sdfg = _make_sdfg(code=code, dtype=dace_dtype)
+    _test_sdfg(sdfg=sdfg, expected=expected, dtype=dtype)
+
+
+@pytest.mark.gpu
+def test_constant_pi_simple():
+    _perform_test(
+            code="math.pi",
+            expected=math.pi
+    )
+
+
+@pytest.mark.gpu
+def test_constant_pi_add():
+    _perform_test(
+            code="-math.pi",
+            expected=-math.pi
+    )
+    _perform_test(
+            code="math.pi + math.pi",
+            expected=2 * math.pi
+    )
+    _perform_test(
+            code="math.pi - math.pi",
+            expected=0.
+    )
+
+
+@pytest.mark.gpu
+def test_constant_pi_mult():
+    _perform_test(
+            code="(math.pi ** 2) * 2",
+            expected=math.pi * math.pi * 2.0
+    )
+    _perform_test(
+            code="math.pi * 2",
+            expected=2 * math.pi
+    )
+    _perform_test(
+            code="math.pi * 2 + math.pi",
+            expected=2 * math.pi + math.pi
+    )
+    _perform_test(
+            code="math.pi * math.pi * 2",
+            expected=math.pi * math.pi * 2.0
+    )
+    _perform_test(
+            code="math.pi / math.pi ",
+            expected=1
+    )
+    _perform_test(
+            code="(math.pi + math.pi) / math.pi ",
+            expected=2
+    )
+    _perform_test(
+            code="(math.pi * math.pi) / math.pi ",
+            expected=math.pi
+    )
+
+
+@pytest.mark.gpu
+def test_constant_pi_fun():
+    _perform_test(
+            code="math.sin(math.pi)",
+            expected=0,
+    )
+    _perform_test(
+            code="math.sin(math.pi * 4)",
+            expected=math.sin(math.pi * 4),
+    )
+    _perform_test(
+            code="math.sin(math.pi * 5)",
+            expected=math.sin(math.pi * 5),
+    )
+    _perform_test(
+            code="math.cos(math.pi * 4)",
+            expected=math.cos(math.pi * 4),
+    )
+    _perform_test(
+            code="math.cos(math.pi * 5)",
+            expected=math.cos(math.pi * 5),
+    )
+    _perform_test(
+            code="math.log(math.pi)",
+            expected=math.log(math.pi),
+    )
+
+
+@pytest.mark.gpu
+def test_constant_nan():
+    _perform_test(
+            code="math.nan",
+            expected=math.nan
+    )
+    _perform_test(
+            code="math.nan + 2",
+            expected=math.nan
+    )
+    _perform_test(
+            code="math.nan + 2.0",
+            expected=math.nan
+    )
+    _perform_test(
+            code="math.sin(math.nan + 2.0)",
+            expected=math.nan
+    )
+    _perform_test(
+            code="math.sin(math.nan + 2.0) ** 2",
+            expected=math.nan
+    )
+
+
+if __name__ == "__main__":
+    test_constant_pi_simple()
+    test_constant_pi_add()
+    test_constant_pi_mult()
+    test_constant_pi_fun()
+    test_constant_nan()