From 2c3e1948407e58c8637605fa73b85c3bf593e374 Mon Sep 17 00:00:00 2001
From: Neil Kichler <neil.kichler@gmail.com>
Date: Tue, 25 Jun 2024 18:02:55 +0200
Subject: [PATCH] put everything into namespace cu:: and fix ADL issues

---
 examples/bisection.cu                         |   1 +
 examples/cuinterval/examples/bisection.cuh    |   8 +-
 include/cuinterval/arithmetic/basic.cuh       |  75 ++++++++-
 include/cuinterval/arithmetic/intrinsic.cuh   |   4 +-
 include/cuinterval/format.h                   |  25 +++
 include/cuinterval/interval.h                 |   6 +-
 src/CMakeLists.txt                            |   1 +
 tests/generated/tests_atan2.cpp               |   2 +-
 tests/generated/tests_c_xsc.cpp               |   2 +-
 tests/generated/tests_filib.cpp               |   2 +-
 .../generated/tests_intervalarithmeticjl.cpp  |   2 +-
 tests/generated/tests_libieeep1788_bool.cpp   |   2 +-
 tests/generated/tests_libieeep1788_cancel.cpp |   2 +-
 tests/generated/tests_libieeep1788_elem.cpp   |   2 +-
 tests/generated/tests_libieeep1788_num.cpp    |   2 +-
 .../generated/tests_libieeep1788_rec_bool.cpp |   2 +-
 tests/generated/tests_libieeep1788_set.cpp    |   2 +-
 tests/generated/tests_mpfi.cpp                |   2 +-
 tests/test_converter.py                       |   4 +-
 tests/tests.h                                 |  16 +-
 tests/tests_bisect.cpp                        |   3 +
 tests/tests_bisect.cu                         |   3 +
 tests/tests_loop.cpp                          |   3 +
 tests/tests_loop.cu                           |   1 +
 tests/tests_ops.cu                            | 140 ++++++++---------
 tests/tests_ops.cuh                           | 148 +++++++++---------
 tests/tests_ops.h                             | 140 ++++++++---------
 tests/tests_utils.h                           |   2 +-
 28 files changed, 347 insertions(+), 255 deletions(-)
 create mode 100644 include/cuinterval/format.h

diff --git a/examples/bisection.cu b/examples/bisection.cu
index fdf62c2..303f191 100644
--- a/examples/bisection.cu
+++ b/examples/bisection.cu
@@ -6,6 +6,7 @@
 
 #include "utils.h"
 
+using cu::interval;
 
 template<typename I>
 __device__ I f(I x)
diff --git a/examples/cuinterval/examples/bisection.cuh b/examples/cuinterval/examples/bisection.cuh
index 6c5dcc2..8bb83cc 100644
--- a/examples/cuinterval/examples/bisection.cuh
+++ b/examples/cuinterval/examples/bisection.cuh
@@ -18,13 +18,13 @@ struct local_stack
     size_type len {};
 };
 
-typedef interval<double> (*fn_t)(interval<double>);
+typedef cu::interval<double> (*fn_t)(cu::interval<double>);
 
 // Example implementation of the bisection method for finding all roots in a given interval.
 template<typename T, int max_depth>
-__global__ void bisection(fn_t f, interval<T> x_init, double tol, interval<T> *roots, std::size_t *max_roots)
+__global__ void bisection(fn_t f, cu::interval<T> x_init, double tol, cu::interval<T> *roots, std::size_t *max_roots)
 {
-    using I = interval<T>;
+    using I = cu::interval<T>;
 
     std::size_t n_roots = 0;
     local_stack<I, max_depth> intervals;
@@ -66,7 +66,7 @@ __global__ void bisection(fn_t f, interval<T> x_init, double tol, interval<T> *r
             }
         } else {
             // interval could still contain a root -> bisect
-            split<T> c = bisect(x, 0.5);
+            cu::split<T> c = bisect(x, 0.5);
             // we do depth-first search which often will not be optimal
             intervals.push(c.upper_half);
             intervals.push(c.lower_half);
diff --git a/include/cuinterval/arithmetic/basic.cuh b/include/cuinterval/arithmetic/basic.cuh
index 4ec1ac7..560746e 100644
--- a/include/cuinterval/arithmetic/basic.cuh
+++ b/include/cuinterval/arithmetic/basic.cuh
@@ -1,12 +1,16 @@
 #ifndef CUINTERVAL_ARITHMETIC_BASIC_CUH
 #define CUINTERVAL_ARITHMETIC_BASIC_CUH
 
-#include <cuinterval/interval.h>
 #include "intrinsic.cuh"
+#include <cuinterval/interval.h>
 
 #include <cassert>
+#include <cmath>
 #include <numbers>
 
+namespace cu
+{
+
 //
 // Constant intervals
 //
@@ -120,6 +124,8 @@ inline constexpr __device__ interval<T> sqrt(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> cbrt(interval<T> x)
 {
+    using std::cbrt;
+
     if (empty(x)) {
         return x;
     }
@@ -144,9 +150,9 @@ inline constexpr __device__ interval<T> recip(interval<T> a)
         } else if (a.lb == zero && zero < a.ub) {
             return { intrinsic::rcp_down(a.ub), intrinsic::pos_inf<T>() };
         } else if (a.lb < zero && zero < a.ub) {
-            return ::entire<T>();
+            return entire<T>();
         } else if (a.lb == zero && zero == a.ub) {
-            return ::empty<T>();
+            return empty<T>();
         }
     }
 
@@ -214,6 +220,8 @@ inline constexpr __device__ interval<T> div(interval<T> x, interval<T> y)
 template<typename T>
 inline constexpr __device__ T mag(interval<T> x)
 {
+    using std::max;
+
     if (empty(x)) {
         return intrinsic::nan<T>();
     }
@@ -223,6 +231,8 @@ inline constexpr __device__ T mag(interval<T> x)
 template<typename T>
 inline constexpr __device__ T mig(interval<T> x)
 {
+    using std::min;
+
     // TODO: we might want to split up the function into the bare interval operation and this part.
     //       we could perhaps use a monad for either result or empty using expected?
     if (empty(x)) {
@@ -264,6 +274,8 @@ inline constexpr __device__ interval<T> abs(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> max(interval<T> x, interval<T> y)
 {
+    using std::max;
+
     if (empty(x) || empty(y)) {
         return empty<T>();
     }
@@ -274,6 +286,8 @@ inline constexpr __device__ interval<T> max(interval<T> x, interval<T> y)
 template<typename T>
 inline constexpr __device__ interval<T> min(interval<T> x, interval<T> y)
 {
+    using std::min;
+
     if (empty(x) || empty(y)) {
         return empty<T>();
     }
@@ -565,6 +579,9 @@ inline constexpr __device__ interval<T> cancel_plus(interval<T> x, interval<T> y
 template<typename T>
 inline constexpr __device__ interval<T> intersection(interval<T> x, interval<T> y)
 {
+    using std::max;
+    using std::min;
+
     // extended
     if (disjoint(x, y)) {
         return empty<T>();
@@ -576,6 +593,9 @@ inline constexpr __device__ interval<T> intersection(interval<T> x, interval<T>
 template<typename T>
 inline constexpr __device__ interval<T> convex_hull(interval<T> x, interval<T> y)
 {
+    using std::max;
+    using std::min;
+
     // extended
     if (empty(x)) {
         return y;
@@ -770,9 +790,9 @@ inline constexpr __device__ interval<T> pown(interval<T> x, std::integral auto n
         return empty<T>();
     }
 
+    using intrinsic::next_after;
     using intrinsic::next_floating;
     using intrinsic::prev_floating;
-    using intrinsic::next_after;
 
     if (n % 2) { // odd power
         if (entire(x)) {
@@ -857,6 +877,8 @@ inline constexpr __device__ interval<T> pow_(interval<T> x, T y)
 template<typename T>
 inline constexpr __device__ interval<T> rootn(interval<T> x, std::integral auto n)
 {
+    using std::pow;
+
     if (empty(x)) {
         return x;
     }
@@ -941,6 +963,10 @@ inline constexpr __device__ unsigned int quadrant_pi(T v)
 template<typename T>
 inline constexpr __device__ interval<T> sin(interval<T> x)
 {
+    using std::max;
+    using std::min;
+    using std::sin;
+
     if (empty(x)) {
         return x;
     }
@@ -1016,6 +1042,10 @@ inline constexpr __device__ interval<T> sin(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> sinpi(interval<T> x)
 {
+    using ::sinpi;
+    using std::max;
+    using std::min;
+
     if (empty(x)) {
         return x;
     }
@@ -1064,6 +1094,10 @@ inline constexpr __device__ interval<T> sinpi(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> cos(interval<T> x)
 {
+    using std::cos;
+    using std::max;
+    using std::min;
+
     if (empty(x)) {
         return x;
     }
@@ -1114,6 +1148,10 @@ inline constexpr __device__ interval<T> cos(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> cospi(interval<T> x)
 {
+    using ::cospi;
+    using std::max;
+    using std::min;
+
     if (empty(x)) {
         return x;
     }
@@ -1161,6 +1199,8 @@ inline constexpr __device__ interval<T> cospi(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> tan(interval<T> x)
 {
+    using std::tan;
+
     if (empty(x)) {
         return x;
     }
@@ -1192,6 +1232,8 @@ inline constexpr __device__ interval<T> tan(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> asin(interval<T> x)
 {
+    using std::asin;
+
     if (empty(x)) {
         return x;
     }
@@ -1207,6 +1249,8 @@ inline constexpr __device__ interval<T> asin(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> acos(interval<T> x)
 {
+    using std::acos;
+
     if (empty(x)) {
         return x;
     }
@@ -1222,6 +1266,8 @@ inline constexpr __device__ interval<T> acos(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> atan(interval<T> x)
 {
+    using std::atan;
+
     if (empty(x)) {
         return x;
     }
@@ -1235,6 +1281,9 @@ inline constexpr __device__ interval<T> atan(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> atan2(interval<T> y, interval<T> x)
 {
+    using std::abs;
+    using std::atan2;
+
     if (empty(x) || empty(y)) {
         return empty<T>();
     }
@@ -1324,6 +1373,8 @@ inline constexpr __device__ interval<T> atan2(interval<T> y, interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> sinh(interval<T> x)
 {
+    using std::sinh;
+
     if (empty(x)) {
         return x;
     }
@@ -1335,6 +1386,8 @@ inline constexpr __device__ interval<T> sinh(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> cosh(interval<T> x)
 {
+    using std::cosh;
+
     if (empty(x)) {
         return x;
     }
@@ -1348,6 +1401,8 @@ inline constexpr __device__ interval<T> cosh(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> tanh(interval<T> x)
 {
+    using std::tanh;
+
     if (empty(x)) {
         return x;
     }
@@ -1361,6 +1416,8 @@ inline constexpr __device__ interval<T> tanh(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> asinh(interval<T> x)
 {
+    using std::asinh;
+
     if (empty(x)) {
         return x;
     }
@@ -1372,6 +1429,8 @@ inline constexpr __device__ interval<T> asinh(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> acosh(interval<T> x)
 {
+    using std::acosh;
+
     if (empty(x)) {
         return x;
     }
@@ -1388,6 +1447,8 @@ inline constexpr __device__ interval<T> acosh(interval<T> x)
 template<typename T>
 inline constexpr __device__ interval<T> atanh(interval<T> x)
 {
+    using std::atanh;
+
     if (empty(x)) {
         return x;
     }
@@ -1465,8 +1526,8 @@ inline constexpr __device__ void mince(interval<T> x, interval<T> *xs, std::size
             xs[i] = empty<T>();
         }
     } else {
-        T lb = x.lb;
-        T ub = x.ub;
+        T lb   = x.lb;
+        T ub   = x.ub;
         T step = (ub - lb) / static_cast<T>(out_xs_size);
 
         for (std::size_t i = 0; i < out_xs_size; i++) {
@@ -1475,4 +1536,6 @@ inline constexpr __device__ void mince(interval<T> x, interval<T> *xs, std::size
     }
 }
 
+} // namespace cu
+
 #endif // CUINTERVAL_ARITHMETIC_BASIC_CUH
diff --git a/include/cuinterval/arithmetic/intrinsic.cuh b/include/cuinterval/arithmetic/intrinsic.cuh
index f9d0157..62ed585 100644
--- a/include/cuinterval/arithmetic/intrinsic.cuh
+++ b/include/cuinterval/arithmetic/intrinsic.cuh
@@ -1,7 +1,7 @@
 #ifndef CUINTERVAL_ARITHMETIC_INTRINSIC_CUH
 #define CUINTERVAL_ARITHMETIC_INTRINSIC_CUH
 
-namespace intrinsic
+namespace cu::intrinsic
 {
 // clang-format off
     template<typename T> inline __device__ T fma_down  (T x, T y, T z);
@@ -104,6 +104,6 @@ namespace intrinsic
     template<> inline __device__ float prev_floating(float x)          { return nextafterf(x, intrinsic::neg_inf<float>()); }
 
 // clang-format on
-} // namespace intrinsic
+} // namespace cu::intrinsic
 
 #endif // CUINTERVAL_ARITHMETIC_INTRINSIC_CUH
diff --git a/include/cuinterval/format.h b/include/cuinterval/format.h
new file mode 100644
index 0000000..66e6044
--- /dev/null
+++ b/include/cuinterval/format.h
@@ -0,0 +1,25 @@
+#ifndef CUINTERVAL_FORMAT_H
+#define CUINTERVAL_FORMAT_H
+
+#include <cuinterval/interval.h>
+
+#include <ostream>
+
+namespace cu
+{
+
+template<typename T>
+std::ostream &operator<<(std::ostream &os, interval<T> x)
+{
+    return os << "[" << x.lb << ", " << x.ub << "]";
+}
+
+template<typename T>
+std::ostream &operator<<(std::ostream &os, split<T> x)
+{
+    return os << "[" << x.lower_half << ", " << x.upper_half << "]";
+}
+
+} // namespace cu
+
+#endif
diff --git a/include/cuinterval/interval.h b/include/cuinterval/interval.h
index d72ad7c..04c4e32 100644
--- a/include/cuinterval/interval.h
+++ b/include/cuinterval/interval.h
@@ -1,6 +1,9 @@
 #ifndef CUINTERVAL_INTERVAL_H
 #define CUINTERVAL_INTERVAL_H
 
+namespace cu
+{
+
 template<typename T>
 struct interval
 {
@@ -11,7 +14,6 @@ struct interval
 template<typename T>
 bool operator==(interval<T> lhs, interval<T> rhs)
 {
-
     auto empty = [](interval<T> x) { return !(x.lb <= x.ub); };
 
     if (empty(lhs) && empty(rhs)) {
@@ -30,4 +32,6 @@ struct split
     auto operator<=>(const split &) const = default;
 };
 
+} // namespace cu
+
 #endif // CUINTERVAL_INTERVAL_H
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 6391082..ced24de 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -5,6 +5,7 @@ set(headers
     "${LIB_PATH}/arithmetic/intrinsic.cuh"
     "${LIB_PATH}/arithmetic/basic.cuh"
     "${LIB_PATH}/cuinterval.h"
+    "${LIB_PATH}/format.h"
 )
 
 add_library(cuinterval "main.cu" ${headers})
diff --git a/tests/generated/tests_atan2.cpp b/tests/generated/tests_atan2.cpp
index 0e1cae2..5fb03b5 100644
--- a/tests/generated/tests_atan2.cpp
+++ b/tests/generated/tests_atan2.cpp
@@ -12,7 +12,7 @@ void tests_atan2(cuda_buffer buffer, cudaStream_t stream, cudaEvent_t event) {
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/generated/tests_c_xsc.cpp b/tests/generated/tests_c_xsc.cpp
index 4759589..b345197 100644
--- a/tests/generated/tests_c_xsc.cpp
+++ b/tests/generated/tests_c_xsc.cpp
@@ -12,7 +12,7 @@ void tests_c_xsc(cuda_buffer buffer, cudaStream_t stream, cudaEvent_t event) {
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/generated/tests_filib.cpp b/tests/generated/tests_filib.cpp
index 2aa8c4c..58a3c29 100644
--- a/tests/generated/tests_filib.cpp
+++ b/tests/generated/tests_filib.cpp
@@ -12,7 +12,7 @@ void tests_filib(cuda_buffer buffer, cudaStream_t stream, cudaEvent_t event) {
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/generated/tests_intervalarithmeticjl.cpp b/tests/generated/tests_intervalarithmeticjl.cpp
index 545abc0..c1ac8ad 100644
--- a/tests/generated/tests_intervalarithmeticjl.cpp
+++ b/tests/generated/tests_intervalarithmeticjl.cpp
@@ -12,7 +12,7 @@ void tests_intervalarithmeticjl(cuda_buffer buffer, cudaStream_t stream, cudaEve
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/generated/tests_libieeep1788_bool.cpp b/tests/generated/tests_libieeep1788_bool.cpp
index 505dea2..02ab2cf 100644
--- a/tests/generated/tests_libieeep1788_bool.cpp
+++ b/tests/generated/tests_libieeep1788_bool.cpp
@@ -12,7 +12,7 @@ void tests_libieeep1788_bool(cuda_buffer buffer, cudaStream_t stream, cudaEvent_
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/generated/tests_libieeep1788_cancel.cpp b/tests/generated/tests_libieeep1788_cancel.cpp
index 892ce66..9c12cf0 100644
--- a/tests/generated/tests_libieeep1788_cancel.cpp
+++ b/tests/generated/tests_libieeep1788_cancel.cpp
@@ -12,7 +12,7 @@ void tests_libieeep1788_cancel(cuda_buffer buffer, cudaStream_t stream, cudaEven
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/generated/tests_libieeep1788_elem.cpp b/tests/generated/tests_libieeep1788_elem.cpp
index db4d291..2e98dbd 100644
--- a/tests/generated/tests_libieeep1788_elem.cpp
+++ b/tests/generated/tests_libieeep1788_elem.cpp
@@ -12,7 +12,7 @@ void tests_libieeep1788_elem(cuda_buffer buffer, cudaStream_t stream, cudaEvent_
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/generated/tests_libieeep1788_num.cpp b/tests/generated/tests_libieeep1788_num.cpp
index ed74f12..4070019 100644
--- a/tests/generated/tests_libieeep1788_num.cpp
+++ b/tests/generated/tests_libieeep1788_num.cpp
@@ -12,7 +12,7 @@ void tests_libieeep1788_num(cuda_buffer buffer, cudaStream_t stream, cudaEvent_t
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/generated/tests_libieeep1788_rec_bool.cpp b/tests/generated/tests_libieeep1788_rec_bool.cpp
index e8697b4..5918e9f 100644
--- a/tests/generated/tests_libieeep1788_rec_bool.cpp
+++ b/tests/generated/tests_libieeep1788_rec_bool.cpp
@@ -12,7 +12,7 @@ void tests_libieeep1788_rec_bool(cuda_buffer buffer, cudaStream_t stream, cudaEv
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/generated/tests_libieeep1788_set.cpp b/tests/generated/tests_libieeep1788_set.cpp
index e85c2c0..127ad6e 100644
--- a/tests/generated/tests_libieeep1788_set.cpp
+++ b/tests/generated/tests_libieeep1788_set.cpp
@@ -12,7 +12,7 @@ void tests_libieeep1788_set(cuda_buffer buffer, cudaStream_t stream, cudaEvent_t
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/generated/tests_mpfi.cpp b/tests/generated/tests_mpfi.cpp
index 385be43..5e6c6d1 100644
--- a/tests/generated/tests_mpfi.cpp
+++ b/tests/generated/tests_mpfi.cpp
@@ -12,7 +12,7 @@ void tests_mpfi(cuda_buffer buffer, cudaStream_t stream, cudaEvent_t event) {
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
diff --git a/tests/test_converter.py b/tests/test_converter.py
index 6e4637b..9766b92 100644
--- a/tests/test_converter.py
+++ b/tests/test_converter.py
@@ -116,7 +116,7 @@ def convert_to_test(file_path):
     using namespace boost::ut;
 
     using T = double;
-    using I = interval<T>;
+    using I = cu::interval<T>;
     using B = bool;
     using N = int;
 
@@ -261,7 +261,7 @@ def generate_kernel_wrappers(ops: dict):
     wrappers_cpp = auto_generated_comment + '#include "tests_ops.cuh"\n#include "tests_common.h"\n'
     wrappers_h = auto_generated_comment + '#include "tests_common.h"\n#include <cuinterval/interval.h>\n'
 
-    template = {I: 'interval<double>',
+    template = {I: 'cu::interval<double>',
                 B: 'bool',
                 T: 'double',
                 N: 'int'}
diff --git a/tests/tests.h b/tests/tests.h
index 0bdab4a..189c3cf 100644
--- a/tests/tests.h
+++ b/tests/tests.h
@@ -16,26 +16,14 @@
 #include <boost/ut.hpp>
 #endif
 
+#include <cuinterval/format.h>
 #include <cuinterval/interval.h>
 
 #include <cmath>
-#include <ostream>
 #include <source_location>
 #include <span>
 #include <vector>
 
-template<typename T>
-std::ostream &operator<<(std::ostream &os, const interval<T> &x)
-{
-    return os << '[' << x.lb << ',' << x.ub << ']';
-}
-
-template<typename T>
-std::ostream &operator<<(std::ostream &os, const split<T> &x)
-{
-    return os << '(' << x.lower_half << ',' << x.upper_half << ')';
-}
-
 template<typename T>
 bool check_within_ulps(T x, T y, std::size_t n, T direction)
 {
@@ -76,7 +64,7 @@ void check_all_equal(T *h_res, std::span<T, N> h_ref, int max_ulps_diff, std::so
         if (h_res[i] != h_res[i] && h_ref[i] != h_ref[i]) // both are NaN
             continue;
 
-        if constexpr (std::is_same_v<T, interval<double>>) {
+        if constexpr (std::is_same_v<T, cu::interval<double>>) {
             if (!empty(h_res[i]) || !empty(h_ref[i])) {
                 bool lb_within_ulps = check_within_ulps(h_res[i].lb, h_ref[i].lb, max_ulps_diff, -std::numeric_limits<double>::infinity());
                 bool ub_within_ulps = check_within_ulps(h_res[i].ub, h_ref[i].ub, max_ulps_diff, std::numeric_limits<double>::infinity());
diff --git a/tests/tests_bisect.cpp b/tests/tests_bisect.cpp
index de07f0d..290b641 100644
--- a/tests/tests_bisect.cpp
+++ b/tests/tests_bisect.cpp
@@ -8,6 +8,9 @@
 #include "tests_ops.h"
 #include "tests_utils.h"
 
+using cu::interval;
+using cu::split;
+
 void test_bisect_call(cudaStream_t stream, int n,
                       interval<double> *x, double *y, split<double> *res);
 
diff --git a/tests/tests_bisect.cu b/tests/tests_bisect.cu
index 8a11adc..d9914b8 100644
--- a/tests/tests_bisect.cu
+++ b/tests/tests_bisect.cu
@@ -9,6 +9,9 @@
 #include "tests_common.h"
 #include "tests_ops.cuh"
 
+using cu::interval;
+using cu::split;
+
 void test_bisect_call(cudaStream_t stream, int n,
                       interval<double> *x, double *y, split<double> *res)
 {
diff --git a/tests/tests_loop.cpp b/tests/tests_loop.cpp
index 77514b4..d7faa03 100644
--- a/tests/tests_loop.cpp
+++ b/tests/tests_loop.cpp
@@ -9,6 +9,9 @@
 #include <omp.h>
 #include <thrust/host_vector.h>
 
+using cu::interval;
+using cu::split;
+
 thrust::host_vector<interval<double>> compute_pi_approximation(cudaStream_t stream);
 
 void tests_pi_approximation(cudaStream_t stream, cudaEvent_t event)
diff --git a/tests/tests_loop.cu b/tests/tests_loop.cu
index 4bd6bce..14cde38 100644
--- a/tests/tests_loop.cu
+++ b/tests/tests_loop.cu
@@ -10,6 +10,7 @@
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/transform.h>
 
+using cu::interval;
 
 struct to_interval_fn
 {
diff --git a/tests/tests_ops.cu b/tests/tests_ops.cu
index 51fdd68..f813bb1 100644
--- a/tests/tests_ops.cu
+++ b/tests/tests_ops.cu
@@ -3,421 +3,421 @@
 #include "tests_common.h"
 
 void tests_pos_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_pos<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_neg_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_neg<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_add_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_add<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_sub_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_sub<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_mul_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_mul<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_div_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_div<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_recip_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_recip<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_sqr_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_sqr<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_sqrt_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_sqrt<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_fma_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_zs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_zs, cu::interval<double> *d_res)
 {
     test_fma<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_zs, d_res); 
 }
             
 void tests_mig_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res)
+                 int n, cu::interval<double> *d_xs, double *d_res)
 {
     test_mig<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_mag_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res)
+                 int n, cu::interval<double> *d_xs, double *d_res)
 {
     test_mag<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_wid_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res)
+                 int n, cu::interval<double> *d_xs, double *d_res)
 {
     test_wid<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_inf_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res)
+                 int n, cu::interval<double> *d_xs, double *d_res)
 {
     test_inf<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_sup_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res)
+                 int n, cu::interval<double> *d_xs, double *d_res)
 {
     test_sup<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_mid_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res)
+                 int n, cu::interval<double> *d_xs, double *d_res)
 {
     test_mid<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_rad_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res)
+                 int n, cu::interval<double> *d_xs, double *d_res)
 {
     test_rad<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_floor_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_floor<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_ceil_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_ceil<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_abs_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_abs<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_min_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_min<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_max_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_max<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_trunc_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_trunc<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_sign_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_sign<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_intersection_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_intersection<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_convexHull_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_convexHull<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_equal_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res)
 {
     test_equal<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_subset_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res)
 {
     test_subset<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_interior_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res)
 {
     test_interior<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_disjoint_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res)
 {
     test_disjoint<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_isEmpty_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, bool *d_res)
+                 int n, cu::interval<double> *d_xs, bool *d_res)
 {
     test_isEmpty<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_isEntire_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, bool *d_res)
+                 int n, cu::interval<double> *d_xs, bool *d_res)
 {
     test_isEntire<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_less_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res)
 {
     test_less<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_strictLess_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res)
 {
     test_strictLess<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_precedes_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res)
 {
     test_precedes<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_strictPrecedes_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res)
 {
     test_strictPrecedes<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_isMember_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, double *d_xs, interval<double> *d_ys, bool *d_res)
+                 int n, double *d_xs, cu::interval<double> *d_ys, bool *d_res)
 {
     test_isMember<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_isSingleton_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, bool *d_res)
+                 int n, cu::interval<double> *d_xs, bool *d_res)
 {
     test_isSingleton<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_isCommonInterval_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, bool *d_res)
+                 int n, cu::interval<double> *d_xs, bool *d_res)
 {
     test_isCommonInterval<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_cancelMinus_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_cancelMinus<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_cancelPlus_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_cancelPlus<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_roundTiesToEven_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_roundTiesToEven<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_roundTiesToAway_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_roundTiesToAway<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_cbrt_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_cbrt<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_exp_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_exp<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_exp2_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_exp2<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_exp10_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_exp10<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_expm1_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_expm1<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_log_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_log<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_log2_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_log2<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_log10_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_log10<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_log1p_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_log1p<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_sin_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_sin<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_cos_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_cos<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_tan_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_tan<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_asin_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_asin<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_acos_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_acos<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_atan_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_atan<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_atan2_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_atan2<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_sinh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_sinh<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_cosh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_cosh<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_tanh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_tanh<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_asinh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_asinh<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_acosh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_acosh<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_atanh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_atanh<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_sinpi_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_sinpi<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_cospi_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res)
 {
     test_cospi<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_res); 
 }
             
 void tests_pown_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, int *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, int *d_ys, cu::interval<double> *d_res)
 {
     test_pown<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_pow_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res)
 {
     test_pow<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
             
 void tests_rootn_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, int *d_ys, interval<double> *d_res)
+                 int n, cu::interval<double> *d_xs, int *d_ys, cu::interval<double> *d_res)
 {
     test_rootn<<<numBlocks, blockSize, 0, stream>>>(n, d_xs, d_ys, d_res); 
 }
diff --git a/tests/tests_ops.cuh b/tests/tests_ops.cuh
index 6a1e1e0..b80bfe2 100644
--- a/tests/tests_ops.cuh
+++ b/tests/tests_ops.cuh
@@ -4,7 +4,7 @@
 #include <cuinterval/cuinterval.h>
 
 template<typename T>
-__global__ void test_pos(int n, interval<T> *x, interval<T> *res)
+__global__ void test_pos(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -13,7 +13,7 @@ __global__ void test_pos(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_neg(int n, interval<T> *x, interval<T> *res)
+__global__ void test_neg(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -22,7 +22,7 @@ __global__ void test_neg(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_recip(int n, interval<T> *x, interval<T> *res)
+__global__ void test_recip(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -31,7 +31,7 @@ __global__ void test_recip(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_sqr(int n, interval<T> *x, interval<T> *res)
+__global__ void test_sqr(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -40,7 +40,7 @@ __global__ void test_sqr(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_sqrt(int n, interval<T> *x, interval<T> *res)
+__global__ void test_sqrt(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -49,7 +49,7 @@ __global__ void test_sqrt(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_cbrt(int n, interval<T> *x, interval<T> *res)
+__global__ void test_cbrt(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -58,7 +58,7 @@ __global__ void test_cbrt(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_add(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_add(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -67,7 +67,7 @@ __global__ void test_add(int n, interval<T> *x, interval<T> *y, interval<T> *res
 }
 
 template<typename T>
-__global__ void test_cancelPlus(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_cancelPlus(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -76,7 +76,7 @@ __global__ void test_cancelPlus(int n, interval<T> *x, interval<T> *y, interval<
 }
 
 template<typename T>
-__global__ void test_cancelMinus(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_cancelMinus(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -85,7 +85,7 @@ __global__ void test_cancelMinus(int n, interval<T> *x, interval<T> *y, interval
 }
 
 template<typename T>
-__global__ void test_sub(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_sub(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -94,7 +94,7 @@ __global__ void test_sub(int n, interval<T> *x, interval<T> *y, interval<T> *res
 }
 
 template<typename T>
-__global__ void test_mul(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_mul(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -103,7 +103,7 @@ __global__ void test_mul(int n, interval<T> *x, interval<T> *y, interval<T> *res
 }
 
 template<typename T>
-__global__ void test_div(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_div(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -112,7 +112,7 @@ __global__ void test_div(int n, interval<T> *x, interval<T> *y, interval<T> *res
 }
 
 template<typename T>
-__global__ void test_fma(int n, interval<T> *x, interval<T> *y, interval<T> *z, interval<T> *res)
+__global__ void test_fma(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *z, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -121,7 +121,7 @@ __global__ void test_fma(int n, interval<T> *x, interval<T> *y, interval<T> *z,
 }
 
 template<typename T>
-__global__ void test_inf(int n, interval<T> *x, T *res)
+__global__ void test_inf(int n, cu::interval<T> *x, T *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -130,7 +130,7 @@ __global__ void test_inf(int n, interval<T> *x, T *res)
 }
 
 template<typename T>
-__global__ void test_sup(int n, interval<T> *x, T *res)
+__global__ void test_sup(int n, cu::interval<T> *x, T *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -139,7 +139,7 @@ __global__ void test_sup(int n, interval<T> *x, T *res)
 }
 
 template<typename T>
-__global__ void test_mid(int n, interval<T> *x, T *res)
+__global__ void test_mid(int n, cu::interval<T> *x, T *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -148,7 +148,7 @@ __global__ void test_mid(int n, interval<T> *x, T *res)
 }
 
 template<typename T>
-__global__ void test_rad(int n, interval<T> *x, T *res)
+__global__ void test_rad(int n, cu::interval<T> *x, T *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -157,7 +157,7 @@ __global__ void test_rad(int n, interval<T> *x, T *res)
 }
 
 template<typename T>
-__global__ void test_mag(int n, interval<T> *x, T *res)
+__global__ void test_mag(int n, cu::interval<T> *x, T *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -166,7 +166,7 @@ __global__ void test_mag(int n, interval<T> *x, T *res)
 }
 
 template<typename T>
-__global__ void test_mig(int n, interval<T> *x, T *res)
+__global__ void test_mig(int n, cu::interval<T> *x, T *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -175,7 +175,7 @@ __global__ void test_mig(int n, interval<T> *x, T *res)
 }
 
 template<typename T>
-__global__ void test_wid(int n, interval<T> *x, T *res)
+__global__ void test_wid(int n, cu::interval<T> *x, T *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -184,7 +184,7 @@ __global__ void test_wid(int n, interval<T> *x, T *res)
 }
 
 template<typename T>
-__global__ void test_floor(int n, interval<T> *x, interval<T> *res)
+__global__ void test_floor(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -193,7 +193,7 @@ __global__ void test_floor(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_ceil(int n, interval<T> *x, interval<T> *res)
+__global__ void test_ceil(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -202,7 +202,7 @@ __global__ void test_ceil(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_trunc(int n, interval<T> *x, interval<T> *res)
+__global__ void test_trunc(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -211,7 +211,7 @@ __global__ void test_trunc(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_sign(int n, interval<T> *x, interval<T> *res)
+__global__ void test_sign(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -220,7 +220,7 @@ __global__ void test_sign(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_abs(int n, interval<T> *x, interval<T> *res)
+__global__ void test_abs(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -229,7 +229,7 @@ __global__ void test_abs(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_min(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_min(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -238,7 +238,7 @@ __global__ void test_min(int n, interval<T> *x, interval<T> *y, interval<T> *res
 }
 
 template<typename T>
-__global__ void test_max(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_max(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -247,7 +247,7 @@ __global__ void test_max(int n, interval<T> *x, interval<T> *y, interval<T> *res
 }
 
 template<typename T>
-__global__ void test_intersection(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_intersection(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -256,7 +256,7 @@ __global__ void test_intersection(int n, interval<T> *x, interval<T> *y, interva
 }
 
 template<typename T>
-__global__ void test_convexHull(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_convexHull(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -265,7 +265,7 @@ __global__ void test_convexHull(int n, interval<T> *x, interval<T> *y, interval<
 }
 
 template<typename T>
-__global__ void test_equal(int n, interval<T> *x, interval<T> *y, bool *res)
+__global__ void test_equal(int n, cu::interval<T> *x, cu::interval<T> *y, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -274,7 +274,7 @@ __global__ void test_equal(int n, interval<T> *x, interval<T> *y, bool *res)
 }
 
 template<typename T>
-__global__ void test_subset(int n, interval<T> *x, interval<T> *y, bool *res)
+__global__ void test_subset(int n, cu::interval<T> *x, cu::interval<T> *y, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -283,7 +283,7 @@ __global__ void test_subset(int n, interval<T> *x, interval<T> *y, bool *res)
 }
 
 template<typename T>
-__global__ void test_interior(int n, interval<T> *x, interval<T> *y, bool *res)
+__global__ void test_interior(int n, cu::interval<T> *x, cu::interval<T> *y, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -292,7 +292,7 @@ __global__ void test_interior(int n, interval<T> *x, interval<T> *y, bool *res)
 }
 
 template<typename T>
-__global__ void test_disjoint(int n, interval<T> *x, interval<T> *y, bool *res)
+__global__ void test_disjoint(int n, cu::interval<T> *x, cu::interval<T> *y, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -301,7 +301,7 @@ __global__ void test_disjoint(int n, interval<T> *x, interval<T> *y, bool *res)
 }
 
 template<typename T>
-__global__ void test_isEmpty(int n, interval<T> *x, bool *res)
+__global__ void test_isEmpty(int n, cu::interval<T> *x, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -310,7 +310,7 @@ __global__ void test_isEmpty(int n, interval<T> *x, bool *res)
 }
 
 template<typename T>
-__global__ void test_isEntire(int n, interval<T> *x, bool *res)
+__global__ void test_isEntire(int n, cu::interval<T> *x, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -319,7 +319,7 @@ __global__ void test_isEntire(int n, interval<T> *x, bool *res)
 }
 
 template<typename T>
-__global__ void test_less(int n, interval<T> *x, interval<T> *y, bool *res)
+__global__ void test_less(int n, cu::interval<T> *x, cu::interval<T> *y, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -328,7 +328,7 @@ __global__ void test_less(int n, interval<T> *x, interval<T> *y, bool *res)
 }
 
 template<typename T>
-__global__ void test_strictLess(int n, interval<T> *x, interval<T> *y, bool *res)
+__global__ void test_strictLess(int n, cu::interval<T> *x, cu::interval<T> *y, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -337,7 +337,7 @@ __global__ void test_strictLess(int n, interval<T> *x, interval<T> *y, bool *res
 }
 
 template<typename T>
-__global__ void test_precedes(int n, interval<T> *x, interval<T> *y, bool *res)
+__global__ void test_precedes(int n, cu::interval<T> *x, cu::interval<T> *y, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -346,7 +346,7 @@ __global__ void test_precedes(int n, interval<T> *x, interval<T> *y, bool *res)
 }
 
 template<typename T>
-__global__ void test_strictPrecedes(int n, interval<T> *x, interval<T> *y, bool *res)
+__global__ void test_strictPrecedes(int n, cu::interval<T> *x, cu::interval<T> *y, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -355,7 +355,7 @@ __global__ void test_strictPrecedes(int n, interval<T> *x, interval<T> *y, bool
 }
 
 template<typename T>
-__global__ void test_isMember(int n, T *x, interval<T> *y, bool *res)
+__global__ void test_isMember(int n, T *x, cu::interval<T> *y, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -364,7 +364,7 @@ __global__ void test_isMember(int n, T *x, interval<T> *y, bool *res)
 }
 
 template<typename T>
-__global__ void test_isSingleton(int n, interval<T> *x, bool *res)
+__global__ void test_isSingleton(int n, cu::interval<T> *x, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -373,7 +373,7 @@ __global__ void test_isSingleton(int n, interval<T> *x, bool *res)
 }
 
 template<typename T>
-__global__ void test_isCommonInterval(int n, interval<T> *x, bool *res)
+__global__ void test_isCommonInterval(int n, cu::interval<T> *x, bool *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -382,7 +382,7 @@ __global__ void test_isCommonInterval(int n, interval<T> *x, bool *res)
 }
 
 template<typename T>
-__global__ void test_roundTiesToEven(int n, interval<T> *x, interval<T> *res)
+__global__ void test_roundTiesToEven(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -391,7 +391,7 @@ __global__ void test_roundTiesToEven(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_roundTiesToAway(int n, interval<T> *x, interval<T> *res)
+__global__ void test_roundTiesToAway(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -400,7 +400,7 @@ __global__ void test_roundTiesToAway(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_exp(int n, interval<T> *x, interval<T> *res)
+__global__ void test_exp(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -409,7 +409,7 @@ __global__ void test_exp(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_exp2(int n, interval<T> *x, interval<T> *res)
+__global__ void test_exp2(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -418,7 +418,7 @@ __global__ void test_exp2(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_exp10(int n, interval<T> *x, interval<T> *res)
+__global__ void test_exp10(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -427,7 +427,7 @@ __global__ void test_exp10(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_expm1(int n, interval<T> *x, interval<T> *res)
+__global__ void test_expm1(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -436,7 +436,7 @@ __global__ void test_expm1(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_log(int n, interval<T> *x, interval<T> *res)
+__global__ void test_log(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -445,7 +445,7 @@ __global__ void test_log(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_log2(int n, interval<T> *x, interval<T> *res)
+__global__ void test_log2(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -454,7 +454,7 @@ __global__ void test_log2(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_log10(int n, interval<T> *x, interval<T> *res)
+__global__ void test_log10(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -463,7 +463,7 @@ __global__ void test_log10(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_log1p(int n, interval<T> *x, interval<T> *res)
+__global__ void test_log1p(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -472,7 +472,7 @@ __global__ void test_log1p(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_pown(int n, interval<T> *x, int p, interval<T> *res)
+__global__ void test_pown(int n, cu::interval<T> *x, int p, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -481,7 +481,7 @@ __global__ void test_pown(int n, interval<T> *x, int p, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_sin(int n, interval<T> *x, interval<T> *res)
+__global__ void test_sin(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -490,7 +490,7 @@ __global__ void test_sin(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_cos(int n, interval<T> *x, interval<T> *res)
+__global__ void test_cos(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -499,7 +499,7 @@ __global__ void test_cos(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_tan(int n, interval<T> *x, interval<T> *res)
+__global__ void test_tan(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -508,7 +508,7 @@ __global__ void test_tan(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_asin(int n, interval<T> *x, interval<T> *res)
+__global__ void test_asin(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -517,7 +517,7 @@ __global__ void test_asin(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_acos(int n, interval<T> *x, interval<T> *res)
+__global__ void test_acos(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -526,7 +526,7 @@ __global__ void test_acos(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_atan(int n, interval<T> *x, interval<T> *res)
+__global__ void test_atan(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -535,7 +535,7 @@ __global__ void test_atan(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_sinh(int n, interval<T> *x, interval<T> *res)
+__global__ void test_sinh(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -544,7 +544,7 @@ __global__ void test_sinh(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_cosh(int n, interval<T> *x, interval<T> *res)
+__global__ void test_cosh(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -553,7 +553,7 @@ __global__ void test_cosh(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_tanh(int n, interval<T> *x, interval<T> *res)
+__global__ void test_tanh(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -562,7 +562,7 @@ __global__ void test_tanh(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_asinh(int n, interval<T> *x, interval<T> *res)
+__global__ void test_asinh(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -571,7 +571,7 @@ __global__ void test_asinh(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_acosh(int n, interval<T> *x, interval<T> *res)
+__global__ void test_acosh(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -580,7 +580,7 @@ __global__ void test_acosh(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_atanh(int n, interval<T> *x, interval<T> *res)
+__global__ void test_atanh(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -589,7 +589,7 @@ __global__ void test_atanh(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_atan2(int n, interval<T> *y, interval<T> *x, interval<T> *res)
+__global__ void test_atan2(int n, cu::interval<T> *y, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -598,7 +598,7 @@ __global__ void test_atan2(int n, interval<T> *y, interval<T> *x, interval<T> *r
 }
 
 template<typename T>
-__global__ void test_sinpi(int n, interval<T> *x, interval<T> *res)
+__global__ void test_sinpi(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -607,7 +607,7 @@ __global__ void test_sinpi(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_cospi(int n, interval<T> *x, interval<T> *res)
+__global__ void test_cospi(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -616,7 +616,7 @@ __global__ void test_cospi(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_cot(int n, interval<T> *x, interval<T> *res)
+__global__ void test_cot(int n, cu::interval<T> *x, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -625,7 +625,7 @@ __global__ void test_cot(int n, interval<T> *x, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_pown(int n, interval<T> *x, int *n_pow, interval<T> *res)
+__global__ void test_pown(int n, cu::interval<T> *x, int *n_pow, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -634,7 +634,7 @@ __global__ void test_pown(int n, interval<T> *x, int *n_pow, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_rootn(int n, interval<T> *x, int *n_pow, interval<T> *res)
+__global__ void test_rootn(int n, cu::interval<T> *x, int *n_pow, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -643,7 +643,7 @@ __global__ void test_rootn(int n, interval<T> *x, int *n_pow, interval<T> *res)
 }
 
 template<typename T>
-__global__ void test_pow(int n, interval<T> *x, interval<T> *y, interval<T> *res)
+__global__ void test_pow(int n, cu::interval<T> *x, cu::interval<T> *y, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -652,7 +652,7 @@ __global__ void test_pow(int n, interval<T> *x, interval<T> *y, interval<T> *res
 }
 
 template<typename T>
-__global__ void test_bisect(int n, interval<T> *x, T *y, split<T> *res)
+__global__ void test_bisect(int n, cu::interval<T> *x, T *y, cu::split<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
@@ -661,7 +661,7 @@ __global__ void test_bisect(int n, interval<T> *x, T *y, split<T> *res)
 }
 
 template<typename T>
-__global__ void test_mince(int n, interval<T> *x, int *d_offsets, interval<T> *res)
+__global__ void test_mince(int n, cu::interval<T> *x, int *d_offsets, cu::interval<T> *res)
 {
     int i = threadIdx.x + blockIdx.x * blockDim.x;
     if (i < n) {
diff --git a/tests/tests_ops.h b/tests/tests_ops.h
index a45b141..5adc6b4 100644
--- a/tests/tests_ops.h
+++ b/tests/tests_ops.h
@@ -3,211 +3,211 @@
 #include <cuinterval/interval.h>
 
 void tests_pos_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_neg_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_add_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_sub_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_mul_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_div_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_recip_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_sqr_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_sqrt_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_fma_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_zs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_zs, cu::interval<double> *d_res);
 
 void tests_mig_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res);
+                 int n, cu::interval<double> *d_xs, double *d_res);
 
 void tests_mag_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res);
+                 int n, cu::interval<double> *d_xs, double *d_res);
 
 void tests_wid_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res);
+                 int n, cu::interval<double> *d_xs, double *d_res);
 
 void tests_inf_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res);
+                 int n, cu::interval<double> *d_xs, double *d_res);
 
 void tests_sup_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res);
+                 int n, cu::interval<double> *d_xs, double *d_res);
 
 void tests_mid_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res);
+                 int n, cu::interval<double> *d_xs, double *d_res);
 
 void tests_rad_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, double *d_res);
+                 int n, cu::interval<double> *d_xs, double *d_res);
 
 void tests_floor_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_ceil_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_abs_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_min_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_max_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_trunc_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_sign_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_intersection_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_convexHull_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_equal_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res);
 
 void tests_subset_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res);
 
 void tests_interior_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res);
 
 void tests_disjoint_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res);
 
 void tests_isEmpty_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, bool *d_res);
+                 int n, cu::interval<double> *d_xs, bool *d_res);
 
 void tests_isEntire_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, bool *d_res);
+                 int n, cu::interval<double> *d_xs, bool *d_res);
 
 void tests_less_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res);
 
 void tests_strictLess_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res);
 
 void tests_precedes_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res);
 
 void tests_strictPrecedes_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, bool *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, bool *d_res);
 
 void tests_isMember_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, double *d_xs, interval<double> *d_ys, bool *d_res);
+                 int n, double *d_xs, cu::interval<double> *d_ys, bool *d_res);
 
 void tests_isSingleton_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, bool *d_res);
+                 int n, cu::interval<double> *d_xs, bool *d_res);
 
 void tests_isCommonInterval_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, bool *d_res);
+                 int n, cu::interval<double> *d_xs, bool *d_res);
 
 void tests_cancelMinus_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_cancelPlus_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_roundTiesToEven_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_roundTiesToAway_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_cbrt_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_exp_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_exp2_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_exp10_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_expm1_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_log_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_log2_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_log10_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_log1p_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_sin_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_cos_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_tan_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_asin_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_acos_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_atan_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_atan2_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_sinh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_cosh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_tanh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_asinh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_acosh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_atanh_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_sinpi_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_cospi_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_res);
 
 void tests_pown_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, int *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, int *d_ys, cu::interval<double> *d_res);
 
 void tests_pow_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, interval<double> *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, cu::interval<double> *d_ys, cu::interval<double> *d_res);
 
 void tests_rootn_call(int numBlocks, int blockSize, cudaStream_t stream,
-                 int n, interval<double> *d_xs, int *d_ys, interval<double> *d_res);
+                 int n, cu::interval<double> *d_xs, int *d_ys, cu::interval<double> *d_res);
diff --git a/tests/tests_utils.h b/tests/tests_utils.h
index 3561230..56ad40b 100644
--- a/tests/tests_utils.h
+++ b/tests/tests_utils.h
@@ -6,7 +6,7 @@
 #include <cassert>
 
 template<typename T>
-void contains(interval<T> x, T y)
+void contains(cu::interval<T> x, T y)
 {
     using namespace boost::ut;