diff --git a/cpp/include/raft/common/span.hpp b/cpp/include/raft/common/span.hpp
index 1e2c36e4fb..7e98d01759 100644
--- a/cpp/include/raft/common/span.hpp
+++ b/cpp/include/raft/common/span.hpp
@@ -21,13 +21,14 @@
 #include <limits>     // numeric_limits
 #include <type_traits>
 
+#include <thrust/functional.h>
 #include <thrust/iterator/reverse_iterator.h>
 
 namespace raft::common {
 
 constexpr std::size_t dynamic_extent = std::numeric_limits<std::size_t>::max();
 
-template <class ElementType, std::size_t Extent>
+template <class ElementType, bool is_device, std::size_t Extent>
 class Span;
 
 namespace detail {
@@ -39,7 +40,7 @@ namespace detail {
  *   - Otherwise, dynamic_extent.
  */
 template <std::size_t Extent, std::size_t Offset, std::size_t Count>
-struct ExtentValue
+struct extent_value_t
   : public std::integral_constant<
       std::size_t, Count != dynamic_extent
                      ? Count
@@ -51,39 +52,28 @@ struct ExtentValue
  * dynamic_extent; otherwise it is std::size_t(sizeof(T)) * N.
  */
 template <typename T, std::size_t Extent>
-struct ExtentAsBytesValue
+struct extent_as_bytes_value_t
   : public std::integral_constant<
       std::size_t, Extent == dynamic_extent ? Extent : sizeof(T) * Extent> {};
 
 template <std::size_t From, std::size_t To>
-struct IsAllowedExtentConversion
+struct is_allowed_extent_conversion_t
   : public std::integral_constant<bool, From == To || From == dynamic_extent ||
                                           To == dynamic_extent> {};
 
 template <class From, class To>
-struct IsAllowedElementTypeConversion
+struct is_allowed_element_type_conversion_t
   : public std::integral_constant<
       bool, std::is_convertible<From (*)[], To (*)[]>::value> {};
 
 template <class T>
-struct IsSpanOracle : std::false_type {};
+struct is_span_oracle_t : std::false_type {};
 
-template <class T, std::size_t Extent>
-struct IsSpanOracle<Span<T, Extent>> : std::true_type {};
+template <class T, bool is_device, std::size_t Extent>
+struct is_span_oracle_t<Span<T, is_device, Extent>> : std::true_type {};
 
 template <class T>
-struct IsSpan : public IsSpanOracle<typename std::remove_cv<T>::type> {};
-
-// Re-implement std algorithms here to adopt CUDA.
-template <typename T>
-struct Less {
-  constexpr bool operator()(const T& _x, const T& _y) const { return _x < _y; }
-};
-
-template <typename T>
-struct Greater {
-  constexpr bool operator()(const T& _x, const T& _y) const { return _x > _y; }
-};
+struct is_span_t : public is_span_oracle_t<typename std::remove_cv<T>::type> {};
 
 template <class InputIt1, class InputIt2, class Compare>
 __host__ __device__ constexpr auto lexicographical_compare(
@@ -105,7 +95,7 @@ __host__ __device__ constexpr auto lexicographical_compare(
  * \brief The span class defined in ISO C++20.  Iterator is defined as plain pointer and
  *        most of the methods have bound check on debug build.
  */
-template <typename T, std::size_t Extent = dynamic_extent>
+template <typename T, bool is_device, std::size_t Extent = dynamic_extent>
 class Span {
  public:
   using element_type = T;
@@ -139,11 +129,12 @@ class Span {
   template <std::size_t N>
   constexpr Span(element_type (&arr)[N]) noexcept : size_(N), data_(&arr[0]) {}
 
-  template <class U, std::size_t OtherExtent,
-            class = typename std::enable_if<
-              detail::IsAllowedElementTypeConversion<U, T>::value &&
-              detail::IsAllowedExtentConversion<OtherExtent, Extent>::value>>
-  constexpr Span(const Span<U, OtherExtent>& _other) noexcept
+  template <
+    class U, std::size_t OtherExtent,
+    class = typename std::enable_if<
+      detail::is_allowed_element_type_conversion_t<U, T>::value &&
+      detail::is_allowed_extent_conversion_t<OtherExtent, Extent>::value>>
+  constexpr Span(const Span<U, is_device, OtherExtent>& _other) noexcept
     : size_(_other.size()), data_(_other.data()) {}
 
   constexpr Span(const Span& _other) noexcept
@@ -208,25 +199,25 @@ class Span {
 
   // Subviews
   template <std::size_t Count>
-  constexpr auto first() const -> Span<element_type, Count> {
+  constexpr auto first() const -> Span<element_type, is_device, Count> {
     assert(Count <= size());
     return {data(), Count};
   }
 
   constexpr auto first(std::size_t _count) const
-    -> Span<element_type, dynamic_extent> {
+    -> Span<element_type, is_device, dynamic_extent> {
     assert(_count <= size());
     return {data(), _count};
   }
 
   template <std::size_t Count>
-  constexpr auto last() const -> Span<element_type, Count> {
+  constexpr auto last() const -> Span<element_type, is_device, Count> {
     assert(Count <= size());
     return {data() + size() - Count, Count};
   }
 
   constexpr auto last(std::size_t _count) const
-    -> Span<element_type, dynamic_extent> {
+    -> Span<element_type, is_device, dynamic_extent> {
     assert(_count <= size());
     return subspan(size() - _count, _count);
   }
@@ -237,7 +228,8 @@ class Span {
    */
   template <std::size_t Offset, std::size_t Count = dynamic_extent>
   constexpr auto subspan() const
-    -> Span<element_type, detail::ExtentValue<Extent, Offset, Count>::value> {
+    -> Span<element_type, is_device,
+            detail::extent_value_t<Extent, Offset, Count>::value> {
     assert((Count == dynamic_extent) ? (Offset <= size())
                                      : (Offset + Count <= size()));
     return {data() + Offset, Count == dynamic_extent ? size() - Offset : Count};
@@ -245,7 +237,7 @@ class Span {
 
   constexpr auto subspan(size_type _offset,
                          size_type _count = dynamic_extent) const
-    -> Span<element_type, dynamic_extent> {
+    -> Span<element_type, is_device, dynamic_extent> {
     assert((_count == dynamic_extent) ? (_offset <= size())
                                       : (_offset + _count <= size()));
     return {data() + _offset,
@@ -257,8 +249,15 @@ class Span {
   pointer data_{nullptr};
 };
 
-template <class T, std::size_t X, class U, std::size_t Y>
-constexpr auto operator==(Span<T, X> l, Span<U, Y> r) -> bool {
+template <typename T, size_t extent = dynamic_extent>
+using host_span = Span<T, false, extent>;
+
+template <typename T, size_t extent = dynamic_extent>
+using device_span = Span<T, true, extent>;
+
+template <class T, std::size_t X, class U, std::size_t Y, bool is_device>
+constexpr auto operator==(Span<T, is_device, X> l, Span<U, is_device, Y> r)
+  -> bool {
   if (l.size() != r.size()) {
     return false;
   }
@@ -271,46 +270,49 @@ constexpr auto operator==(Span<T, X> l, Span<U, Y> r) -> bool {
   return true;
 }
 
-template <class T, std::size_t X, class U, std::size_t Y>
-constexpr auto operator!=(Span<T, X> l, Span<U, Y> r) {
+template <class T, std::size_t X, class U, std::size_t Y, bool is_device>
+constexpr auto operator!=(Span<T, is_device, X> l, Span<U, is_device, Y> r) {
   return !(l == r);
 }
 
-template <class T, std::size_t X, class U, std::size_t Y>
-constexpr auto operator<(Span<T, X> l, Span<U, Y> r) {
+template <class T, std::size_t X, class U, std::size_t Y, bool is_device>
+constexpr auto operator<(Span<T, is_device, X> l, Span<U, is_device, Y> r) {
   return detail::lexicographical_compare<
-    typename Span<T, X>::iterator, typename Span<U, Y>::iterator,
-    detail::Less<typename Span<T, X>::element_type>>(l.begin(), l.end(),
-                                                     r.begin(), r.end());
+    typename Span<T, is_device, X>::iterator,
+    typename Span<U, is_device, Y>::iterator,
+    thrust::less<typename Span<T, is_device, X>::element_type>>(
+    l.begin(), l.end(), r.begin(), r.end());
 }
 
-template <class T, std::size_t X, class U, std::size_t Y>
-constexpr auto operator<=(Span<T, X> l, Span<U, Y> r) {
+template <class T, std::size_t X, class U, std::size_t Y, bool is_device>
+constexpr auto operator<=(Span<T, is_device, X> l, Span<U, is_device, Y> r) {
   return !(l > r);
 }
 
-template <class T, std::size_t X, class U, std::size_t Y>
-constexpr auto operator>(Span<T, X> l, Span<U, Y> r) {
+template <class T, std::size_t X, class U, std::size_t Y, bool is_device>
+constexpr auto operator>(Span<T, is_device, X> l, Span<U, is_device, Y> r) {
   return detail::lexicographical_compare<
-    typename Span<T, X>::iterator, typename Span<U, Y>::iterator,
-    detail::Greater<typename Span<T, X>::element_type>>(l.begin(), l.end(),
-                                                        r.begin(), r.end());
+    typename Span<T, is_device, X>::iterator,
+    typename Span<U, is_device, Y>::iterator,
+    thrust::greater<typename Span<T, is_device, X>::element_type>>(
+    l.begin(), l.end(), r.begin(), r.end());
 }
 
-template <class T, std::size_t X, class U, std::size_t Y>
-constexpr auto operator>=(Span<T, X> l, Span<U, Y> r) {
+template <class T, std::size_t X, class U, std::size_t Y, bool is_device>
+constexpr auto operator>=(Span<T, is_device, X> l, Span<U, is_device, Y> r) {
   return !(l < r);
 }
 
-template <class T, std::size_t E>
-auto as_bytes(Span<T, E> s) noexcept
-  -> Span<const std::byte, detail::ExtentAsBytesValue<T, E>::value> {
+template <class T, bool is_device, std::size_t E>
+auto as_bytes(Span<T, is_device, E> s) noexcept
+  -> Span<const std::byte, is_device,
+          detail::extent_as_bytes_value_t<T, E>::value> {
   return {reinterpret_cast<const std::byte*>(s.data()), s.size_bytes()};
 }
 
-template <class T, std::size_t E>
-auto as_writable_bytes(Span<T, E> s) noexcept
-  -> Span<std::byte, detail::ExtentAsBytesValue<T, E>::value> {
+template <class T, bool is_device, std::size_t E>
+auto as_writable_bytes(Span<T, is_device, E> s) noexcept
+  -> Span<std::byte, is_device, detail::extent_as_bytes_value_t<T, E>::value> {
   return {reinterpret_cast<std::byte*>(s.data()), s.size_bytes()};
 }
 }  // namespace raft::common
diff --git a/cpp/test/common/span.cpp b/cpp/test/common/span.cpp
index 700eb20b1d..7659f7462b 100644
--- a/cpp/test/common/span.cpp
+++ b/cpp/test/common/span.cpp
@@ -1,39 +1,54 @@
+/*
+ * Copyright (c) 2021, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
 #include <gtest/gtest.h>
 #include <numeric>  // iota
 #include <raft/common/span.hpp>
-#include "test_span.h"
+#include "test_span.hpp"
 
 namespace raft::common {
 TEST(Span, DlfConstructors) {
   // Dynamic extent
   {
-    Span<int> s;
+    host_span<int> s;
     ASSERT_EQ(s.size(), 0);
     ASSERT_EQ(s.data(), nullptr);
 
-    Span<int const> cs;
+    host_span<int const> cs;
     ASSERT_EQ(cs.size(), 0);
     ASSERT_EQ(cs.data(), nullptr);
   }
 
   // Static extent
   {
-    Span<int, 0> s;
+    host_span<int, 0> s;
     ASSERT_EQ(s.size(), 0);
     ASSERT_EQ(s.data(), nullptr);
 
-    Span<int const, 0> cs;
+    host_span<int const, 0> cs;
     ASSERT_EQ(cs.size(), 0);
     ASSERT_EQ(cs.data(), nullptr);
   }
 
   // Init list.
   {
-    Span<float> s{};
+    host_span<float> s{};
     ASSERT_EQ(s.size(), 0);
     ASSERT_EQ(s.data(), nullptr);
 
-    Span<int const> cs{};
+    host_span<int const> cs{};
     ASSERT_EQ(cs.size(), 0);
     ASSERT_EQ(cs.data(), nullptr);
   }
@@ -42,21 +57,23 @@ TEST(Span, DlfConstructors) {
 TEST(Span, FromNullPtr) {
   // dynamic extent
   {
-    Span<float> s{nullptr, static_cast<Span<float>::size_type>(0)};
+    host_span<float> s{nullptr, static_cast<host_span<float>::size_type>(0)};
     ASSERT_EQ(s.size(), 0);
     ASSERT_EQ(s.data(), nullptr);
 
-    Span<float const> cs{nullptr, static_cast<Span<float>::size_type>(0)};
+    host_span<float const> cs{nullptr,
+                              static_cast<host_span<float>::size_type>(0)};
     ASSERT_EQ(cs.size(), 0);
     ASSERT_EQ(cs.data(), nullptr);
   }
   // static extent
   {
-    Span<float, 0> s{nullptr, static_cast<Span<float>::size_type>(0)};
+    host_span<float, 0> s{nullptr, static_cast<host_span<float>::size_type>(0)};
     ASSERT_EQ(s.size(), 0);
     ASSERT_EQ(s.data(), nullptr);
 
-    Span<float const, 0> cs{nullptr, static_cast<Span<float>::size_type>(0)};
+    host_span<float const, 0> cs{nullptr,
+                                 static_cast<host_span<float>::size_type>(0)};
     ASSERT_EQ(cs.size(), 0);
     ASSERT_EQ(cs.data(), nullptr);
   }
@@ -68,26 +85,26 @@ TEST(Span, FromPtrLen) {
 
   // static extent
   {
-    Span<float> s(arr, 16);
+    host_span<float> s(arr, 16);
     ASSERT_EQ(s.size(), 16);
     ASSERT_EQ(s.data(), arr);
 
-    for (Span<float>::size_type i = 0; i < 16; ++i) {
+    for (host_span<float>::size_type i = 0; i < 16; ++i) {
       ASSERT_EQ(s[i], arr[i]);
     }
 
-    Span<float const> cs(arr, 16);
+    host_span<float const> cs(arr, 16);
     ASSERT_EQ(cs.size(), 16);
     ASSERT_EQ(cs.data(), arr);
 
-    for (Span<float const>::size_type i = 0; i < 16; ++i) {
+    for (host_span<float const>::size_type i = 0; i < 16; ++i) {
       ASSERT_EQ(cs[i], arr[i]);
     }
   }
 
   // dynamic extent
   {
-    Span<float, 16> s(arr, 16);
+    host_span<float, 16> s(arr, 16);
     ASSERT_EQ(s.size(), 16);
     ASSERT_EQ(s.data(), arr);
 
@@ -95,11 +112,11 @@ TEST(Span, FromPtrLen) {
       ASSERT_EQ(s[i], arr[i]);
     }
 
-    Span<float const, 16> cs(arr, 16);
+    host_span<float const, 16> cs(arr, 16);
     ASSERT_EQ(cs.size(), 16);
     ASSERT_EQ(cs.data(), arr);
 
-    for (Span<float const>::size_type i = 0; i < 16; ++i) {
+    for (host_span<float const>::size_type i = 0; i < 16; ++i) {
       ASSERT_EQ(cs[i], arr[i]);
     }
   }
@@ -107,47 +124,47 @@ TEST(Span, FromPtrLen) {
 
 TEST(Span, FromFirstLast) {
   float arr[16];
-  InitializeRange(arr, arr+16);
+  InitializeRange(arr, arr + 16);
 
   // dynamic extent
   {
-    Span<float> s (arr, arr + 16);
-    ASSERT_EQ (s.size(), 16);
-    ASSERT_EQ (s.data(), arr);
-    ASSERT_EQ (s.data() + s.size(), arr + 16);
+    host_span<float> s(arr, arr + 16);
+    ASSERT_EQ(s.size(), 16);
+    ASSERT_EQ(s.data(), arr);
+    ASSERT_EQ(s.data() + s.size(), arr + 16);
 
     for (size_t i = 0; i < 16; ++i) {
-      ASSERT_EQ (s[i], arr[i]);
+      ASSERT_EQ(s[i], arr[i]);
     }
 
-    Span<float const> cs (arr, arr + 16);
-    ASSERT_EQ (cs.size(), 16);
-    ASSERT_EQ (cs.data(), arr);
-    ASSERT_EQ (cs.data() + cs.size(), arr + 16);
+    host_span<float const> cs(arr, arr + 16);
+    ASSERT_EQ(cs.size(), 16);
+    ASSERT_EQ(cs.data(), arr);
+    ASSERT_EQ(cs.data() + cs.size(), arr + 16);
 
     for (size_t i = 0; i < 16; ++i) {
-      ASSERT_EQ (cs[i], arr[i]);
+      ASSERT_EQ(cs[i], arr[i]);
     }
   }
 
   // static extent
   {
-    Span<float, 16> s (arr, arr + 16);
-    ASSERT_EQ (s.size(), 16);
-    ASSERT_EQ (s.data(), arr);
-    ASSERT_EQ (s.data() + s.size(), arr + 16);
+    host_span<float, 16> s(arr, arr + 16);
+    ASSERT_EQ(s.size(), 16);
+    ASSERT_EQ(s.data(), arr);
+    ASSERT_EQ(s.data() + s.size(), arr + 16);
 
     for (size_t i = 0; i < 16; ++i) {
-      ASSERT_EQ (s[i], arr[i]);
+      ASSERT_EQ(s[i], arr[i]);
     }
 
-    Span<float const> cs (arr, arr + 16);
-    ASSERT_EQ (cs.size(), 16);
-    ASSERT_EQ (cs.data(), arr);
-    ASSERT_EQ (cs.data() + cs.size(), arr + 16);
+    host_span<float const> cs(arr, arr + 16);
+    ASSERT_EQ(cs.size(), 16);
+    ASSERT_EQ(cs.data(), arr);
+    ASSERT_EQ(cs.data() + cs.size(), arr + 16);
 
     for (size_t i = 0; i < 16; ++i) {
-      ASSERT_EQ (cs[i], arr[i]);
+      ASSERT_EQ(cs[i], arr[i]);
     }
   }
 }
@@ -160,11 +177,10 @@ struct DerivedClass : public BaseClass {
 };
 
 TEST(Span, FromOther) {
-
   // convert constructor
   {
-    Span<DerivedClass> derived;
-    Span<BaseClass> base { derived };
+    host_span<DerivedClass> derived;
+    host_span<BaseClass> base{derived};
     ASSERT_EQ(base.size(), derived.size());
     ASSERT_EQ(base.data(), derived.data());
   }
@@ -174,8 +190,8 @@ TEST(Span, FromOther) {
 
   // default copy constructor
   {
-    Span<float> s0 (arr);
-    Span<float> s1 (s0);
+    host_span<float> s0(arr);
+    host_span<float> s1(s0);
     ASSERT_EQ(s0.size(), s1.size());
     ASSERT_EQ(s0.data(), s1.data());
   }
@@ -186,7 +202,7 @@ TEST(Span, FromArray) {
   InitializeRange(arr, arr + 16);
 
   {
-    Span<float> s (arr);
+    host_span<float> s(arr);
     ASSERT_EQ(&arr[0], s.data());
     ASSERT_EQ(s.size(), 16);
     for (size_t i = 0; i < 16; ++i) {
@@ -195,7 +211,7 @@ TEST(Span, FromArray) {
   }
 
   {
-    Span<float, 16> s (arr);
+    host_span<float, 16> s(arr);
     ASSERT_EQ(&arr[0], s.data());
     ASSERT_EQ(s.size(), 16);
     for (size_t i = 0; i < 16; ++i) {
@@ -206,13 +222,13 @@ TEST(Span, FromArray) {
 
 TEST(Span, Assignment) {
   int status = 1;
-  TestAssignment{&status}();
+  TestAssignment<false>{&status}();
   ASSERT_EQ(status, 1);
 }
 
 TEST(Span, BeginEnd) {
   int status = 1;
-  TestBeginEnd{&status}();
+  TestBeginEnd<false>{&status}();
   ASSERT_EQ(status, 1);
 }
 
@@ -220,7 +236,7 @@ TEST(Span, ElementAccess) {
   float arr[16];
   InitializeRange(arr, arr + 16);
 
-  Span<float> s (arr);
+  host_span<float> s(arr);
   size_t j = 0;
   for (auto i : s) {
     ASSERT_EQ(i, arr[j]);
@@ -228,49 +244,35 @@ TEST(Span, ElementAccess) {
   }
 }
 
-
 TEST(Span, Obversers) {
   int status = 1;
-  TestObservers{&status}();
+  TestObservers<false>{&status}();
   ASSERT_EQ(status, 1);
 }
 
 TEST(Span, FrontBack) {
   {
     float arr[4]{0, 1, 2, 3};
-    Span<float, 4> s(arr);
+    host_span<float, 4> s(arr);
     ASSERT_EQ(s.front(), 0);
     ASSERT_EQ(s.back(), 3);
   }
   {
     std::vector<double> arr{0, 1, 2, 3};
-    Span<double> s(arr.data(), arr.size());
+    host_span<double> s(arr.data(), arr.size());
     ASSERT_EQ(s.front(), 0);
     ASSERT_EQ(s.back(), 3);
   }
 }
 
-TEST(SpanDeathTest, FrontBack) {
-  {
-    Span<float, 0> s;
-    EXPECT_DEATH(s.front(), "");
-    EXPECT_DEATH(s.back(), "");
-  }
-  {
-    Span<float> s;
-    EXPECT_DEATH(s.front(), "");
-    EXPECT_DEATH(s.back(), "");
-  }
-}
-
 TEST(Span, FirstLast) {
   // static extent
   {
     float arr[16];
     InitializeRange(arr, arr + 16);
 
-    Span<float> s (arr);
-    Span<float, 4> first = s.first<4>();
+    host_span<float> s(arr);
+    host_span<float, 4> first = s.first<4>();
 
     ASSERT_EQ(first.size(), 4);
     ASSERT_EQ(first.data(), arr);
@@ -284,14 +286,14 @@ TEST(Span, FirstLast) {
     float arr[16];
     InitializeRange(arr, arr + 16);
 
-    Span<float> s (arr);
-    Span<float, 4> last = s.last<4>();
+    host_span<float> s(arr);
+    host_span<float, 4> last = s.last<4>();
 
     ASSERT_EQ(last.size(), 4);
     ASSERT_EQ(last.data(), arr + 12);
 
     for (size_t i = 0; i < last.size(); ++i) {
-      ASSERT_EQ(last[i], arr[i+12]);
+      ASSERT_EQ(last[i], arr[i + 12]);
     }
   }
 
@@ -299,8 +301,8 @@ TEST(Span, FirstLast) {
   {
     float *arr = new float[16];
     InitializeRange(arr, arr + 16);
-    Span<float> s (arr, 16);
-    Span<float> first = s.first(4);
+    host_span<float> s(arr, 16);
+    host_span<float> first = s.first(4);
 
     ASSERT_EQ(first.size(), 4);
     ASSERT_EQ(first.data(), s.data());
@@ -309,14 +311,14 @@ TEST(Span, FirstLast) {
       ASSERT_EQ(first[i], s[i]);
     }
 
-    delete [] arr;
+    delete[] arr;
   }
 
   {
     float *arr = new float[16];
     InitializeRange(arr, arr + 16);
-    Span<float> s (arr, 16);
-    Span<float> last = s.last(4);
+    host_span<float> s(arr, 16);
+    host_span<float> last = s.last(4);
 
     ASSERT_EQ(last.size(), 4);
     ASSERT_EQ(last.data(), s.data() + 12);
@@ -325,13 +327,13 @@ TEST(Span, FirstLast) {
       ASSERT_EQ(s[12 + i], last[i]);
     }
 
-    delete [] arr;
+    delete[] arr;
   }
 }
 
 TEST(Span, Subspan) {
-  int arr[16] {0};
-  Span<int> s1 (arr);
+  int arr[16]{0};
+  host_span<int> s1(arr);
   auto s2 = s1.subspan<4>();
   ASSERT_EQ(s1.size() - 4, s2.size());
 
@@ -346,25 +348,25 @@ TEST(Span, Subspan) {
 
 TEST(Span, Compare) {
   int status = 1;
-  TestCompare{&status}();
+  TestCompare<false>{&status}();
   ASSERT_EQ(status, 1);
 }
 
 TEST(Span, AsBytes) {
   int status = 1;
-  TestAsBytes{&status}();
+  TestAsBytes<false>{&status}();
   ASSERT_EQ(status, 1);
 }
 
 TEST(Span, AsWritableBytes) {
   int status = 1;
-  TestAsWritableBytes{&status}();
+  TestAsWritableBytes<false>{&status}();
   ASSERT_EQ(status, 1);
 }
 
 TEST(Span, Empty) {
   {
-    Span<float> s {nullptr, static_cast<Span<float>::size_type>(0)};
+    host_span<float> s{nullptr, static_cast<host_span<float>::size_type>(0)};
     auto res = s.subspan(0);
     ASSERT_EQ(res.data(), nullptr);
     ASSERT_EQ(res.size(), 0);
@@ -375,7 +377,7 @@ TEST(Span, Empty) {
   }
 
   {
-    Span<float, 0> s {nullptr, static_cast<Span<float>::size_type>(0)};
+    host_span<float, 0> s{nullptr, static_cast<host_span<float>::size_type>(0)};
     auto res = s.subspan(0);
     ASSERT_EQ(res.data(), nullptr);
     ASSERT_EQ(res.size(), 0);
diff --git a/cpp/test/common/span.cu b/cpp/test/common/span.cu
index 108a7ee583..3356fe40f4 100644
--- a/cpp/test/common/span.cu
+++ b/cpp/test/common/span.cu
@@ -1,3 +1,18 @@
+/*
+ * Copyright (c) 2021, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
 #include <gtest/gtest.h>
 #include <raft/cudart_utils.h>
 #include <thrust/device_vector.h>
@@ -5,7 +20,7 @@
 #include <numeric>  // iota
 #include <raft/common/span.hpp>
 #include <raft/cuda_utils.cuh>
-#include "test_span.h"
+#include "test_span.hpp"
 
 namespace raft::common {
 struct TestStatus {
@@ -31,7 +46,7 @@ struct TestStatus {
   int* Data() { return status_; }
 };
 
-__global__ void TestFromOtherKernel(Span<float> span) {
+__global__ void TestFromOtherKernel(device_span<float> span) {
   // don't get optimized out
   size_t idx = threadIdx.x + blockIdx.x * blockDim.x;
 
@@ -40,7 +55,7 @@ __global__ void TestFromOtherKernel(Span<float> span) {
   }
 }
 // Test converting different T
-__global__ void TestFromOtherKernelConst(Span<float const, 16> span) {
+__global__ void TestFromOtherKernelConst(device_span<float const, 16> span) {
   // don't get optimized out
   size_t idx = threadIdx.x + blockIdx.x * blockDim.x;
 
@@ -60,20 +75,20 @@ TEST(GPUSpan, FromOther) {
   thrust::copy(h_vec.begin(), h_vec.end(), d_vec.begin());
   // dynamic extent
   {
-    Span<float> span(d_vec.data().get(), d_vec.size());
+    device_span<float> span(d_vec.data().get(), d_vec.size());
     TestFromOtherKernel<<<1, 16>>>(span);
   }
   {
-    Span<float> span(d_vec.data().get(), d_vec.size());
+    device_span<float> span(d_vec.data().get(), d_vec.size());
     TestFromOtherKernelConst<<<1, 16>>>(span);
   }
   // static extent
   {
-    Span<float, 16> span(d_vec.data().get(), d_vec.data().get() + 16);
+    device_span<float, 16> span(d_vec.data().get(), d_vec.data().get() + 16);
     TestFromOtherKernel<<<1, 16>>>(span);
   }
   {
-    Span<float, 16> span(d_vec.data().get(), d_vec.data().get() + 16);
+    device_span<float, 16> span(d_vec.data().get(), d_vec.data().get() + 16);
     TestFromOtherKernelConst<<<1, 16>>>(span);
   }
 }
@@ -82,7 +97,7 @@ TEST(GPUSpan, Assignment) {
   CUDA_CHECK(cudaSetDevice(0));
   TestStatus status;
   thrust::for_each_n(thrust::make_counting_iterator(0ul), 16,
-                     TestAssignment{status.Data()});
+                     TestAssignment<true>{status.Data()});
   ASSERT_EQ(status.Get(), 1);
 }
 
@@ -123,7 +138,7 @@ TEST(GPUSpan, WithTrust) {
   // Can't initialize span with device_vector, since d_vec.data() is not raw
   // pointer
   {
-    Span<float> s(d_vec.data().get(), d_vec.size());
+    device_span<float> s(d_vec.data().get(), d_vec.size());
 
     ASSERT_EQ(d_vec.size(), s.size());
     ASSERT_EQ(d_vec.data().get(), s.data());
@@ -133,7 +148,7 @@ TEST(GPUSpan, WithTrust) {
     TestStatus status;
     thrust::device_vector<float> d_vec1(d_vec.size());
     thrust::copy(thrust::device, d_vec.begin(), d_vec.end(), d_vec1.begin());
-    Span<float> s(d_vec1.data().get(), d_vec.size());
+    device_span<float> s(d_vec1.data().get(), d_vec.size());
 
     thrust::for_each_n(thrust::make_counting_iterator(0ul), 16,
                        TestEqual<float>{thrust::raw_pointer_cast(d_vec1.data()),
@@ -146,7 +161,7 @@ TEST(GPUSpan, BeginEnd) {
   CUDA_CHECK(cudaSetDevice(0));
   TestStatus status;
   thrust::for_each_n(thrust::make_counting_iterator(0ul), 16,
-                     TestBeginEnd{status.Data()});
+                     TestBeginEnd<true>{status.Data()});
   ASSERT_EQ(status.Get(), 1);
 }
 
@@ -154,11 +169,11 @@ TEST(GPUSpan, RBeginREnd) {
   CUDA_CHECK(cudaSetDevice(0));
   TestStatus status;
   thrust::for_each_n(thrust::make_counting_iterator(0ul), 16,
-                     TestRBeginREnd{status.Data()});
+                     TestRBeginREnd<true>{status.Data()});
   ASSERT_EQ(status.Get(), 1);
 }
 
-__global__ void TestModifyKernel(Span<float> span) {
+__global__ void TestModifyKernel(device_span<float> span) {
   size_t idx = threadIdx.x + blockIdx.x * blockDim.x;
 
   if (idx >= span.size()) {
@@ -174,7 +189,7 @@ TEST(GPUSpan, Modify) {
   thrust::device_vector<float> d_vec(h_vec.size());
   thrust::copy(h_vec.begin(), h_vec.end(), d_vec.begin());
 
-  Span<float> span(d_vec.data().get(), d_vec.size());
+  device_span<float> span(d_vec.data().get(), d_vec.size());
 
   TestModifyKernel<<<1, 16>>>(span);
 
@@ -187,16 +202,16 @@ TEST(GPUSpan, Observers) {
   CUDA_CHECK(cudaSetDevice(0));
   TestStatus status;
   thrust::for_each_n(thrust::make_counting_iterator(0ul), 16,
-                     TestObservers{status.Data()});
+                     TestObservers<true>{status.Data()});
   ASSERT_EQ(status.Get(), 1);
 }
 
 struct TestElementAccess {
  private:
-  Span<float> span_;
+  device_span<float> span_;
 
  public:
-  HD explicit TestElementAccess(Span<float> _span) : span_(_span) {}
+  HD explicit TestElementAccess(device_span<float> _span) : span_(_span) {}
 
   HD float operator()(size_t _idx) {
     float tmp = span_[_idx];
diff --git a/cpp/test/common/test_span.h b/cpp/test/common/test_span.hpp
similarity index 67%
rename from cpp/test/common/test_span.h
rename to cpp/test/common/test_span.hpp
index 6ec377085d..0a03e378b7 100644
--- a/cpp/test/common/test_span.h
+++ b/cpp/test/common/test_span.hpp
@@ -1,3 +1,19 @@
+/*
+ * Copyright (c) 2021, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#pragma once
 #include <raft/common/span.hpp>
 
 namespace raft::common {
@@ -30,6 +46,7 @@ struct TestTestStatus {
   }
 };
 
+template <bool is_device>
 struct TestAssignment {
   int* status_;
 
@@ -37,11 +54,11 @@ struct TestAssignment {
 
   __host__ __device__ void operator()() { this->operator()(0); }
   __host__ __device__ void operator()(int _idx) {
-    Span<float> s1;
+    Span<float, is_device> s1;
 
     float arr[] = {3, 4, 5};
 
-    Span<const float> s2 = arr;
+    Span<const float, is_device> s2 = arr;
     SPAN_ASSERT_TRUE(s2.size() == 3, status_);
     SPAN_ASSERT_TRUE(s2.data() == &arr[0], status_);
 
@@ -50,6 +67,7 @@ struct TestAssignment {
   }
 };
 
+template <bool is_device>
 struct TestBeginEnd {
   int* status_;
 
@@ -60,9 +78,9 @@ struct TestBeginEnd {
     float arr[16];
     InitializeRange(arr, arr + 16);
 
-    Span<float> s(arr);
-    Span<float>::iterator beg{s.begin()};
-    Span<float>::iterator end{s.end()};
+    Span<float, is_device> s(arr);
+    typename Span<float, is_device>::iterator beg{s.begin()};
+    typename Span<float, is_device>::iterator end{s.end()};
 
     SPAN_ASSERT_TRUE(end == beg + 16, status_);
     SPAN_ASSERT_TRUE(*beg == arr[0], status_);
@@ -70,6 +88,7 @@ struct TestBeginEnd {
   }
 };
 
+template <bool is_device>
 struct TestRBeginREnd {
   int* status_;
 
@@ -80,10 +99,10 @@ struct TestRBeginREnd {
     float arr[16];
     InitializeRange(arr, arr + 16);
 
-    Span<float> s(arr);
+    Span<float, is_device> s(arr);
     s.rbegin();
-    Span<float>::reverse_iterator rbeg{s.rbegin()};
-    Span<float>::reverse_iterator rend{s.rend()};
+    typename Span<float, is_device>::reverse_iterator rbeg{s.rbegin()};
+    typename Span<float, is_device>::reverse_iterator rend{s.rend()};
 
     SPAN_ASSERT_TRUE(rbeg + 16 == rend, status_);
     SPAN_ASSERT_TRUE(*(rbeg) == arr[15], status_);
@@ -91,6 +110,7 @@ struct TestRBeginREnd {
   }
 };
 
+template <bool is_device>
 struct TestObservers {
   int* status_;
 
@@ -101,14 +121,15 @@ struct TestObservers {
     // empty
     {
       float* arr = nullptr;
-      Span<float> s(arr, static_cast<Span<float>::size_type>(0));
+      Span<float, is_device> s(
+        arr, static_cast<typename Span<float, is_device>::size_type>(0));
       SPAN_ASSERT_TRUE(s.empty(), status_);
     }
 
     // size, size_types
     {
       float* arr = new float[16];
-      Span<float> s(arr, 16);
+      Span<float, is_device> s(arr, 16);
       SPAN_ASSERT_TRUE(s.size() == 16, status_);
       SPAN_ASSERT_TRUE(s.size_bytes() == 16 * sizeof(float), status_);
       delete[] arr;
@@ -116,6 +137,7 @@ struct TestObservers {
   }
 };
 
+template <bool is_device>
 struct TestCompare {
   int* status_;
 
@@ -127,8 +149,8 @@ struct TestCompare {
     InitializeRange(lhs_arr, lhs_arr + 16);
     InitializeRange(rhs_arr, rhs_arr + 16);
 
-    Span<float> lhs(lhs_arr);
-    Span<float> rhs(rhs_arr);
+    Span<float, is_device> lhs(lhs_arr);
+    Span<float, is_device> rhs(rhs_arr);
 
     SPAN_ASSERT_TRUE(lhs == rhs, status_);
     SPAN_ASSERT_FALSE(lhs != rhs, status_);
@@ -144,6 +166,7 @@ struct TestCompare {
   }
 };
 
+template <bool is_device>
 struct TestAsBytes {
   int* status_;
 
@@ -155,8 +178,8 @@ struct TestAsBytes {
     InitializeRange(arr, arr + 16);
 
     {
-      const Span<const float> s{arr};
-      const Span<const std::byte> bs = as_bytes(s);
+      const Span<const float, is_device> s{arr};
+      const Span<const std::byte, is_device> bs = as_bytes(s);
       SPAN_ASSERT_TRUE(bs.size() == s.size_bytes(), status_);
       SPAN_ASSERT_TRUE(static_cast<const void*>(bs.data()) ==
                          static_cast<const void*>(s.data()),
@@ -164,8 +187,8 @@ struct TestAsBytes {
     }
 
     {
-      Span<float> s;
-      const Span<const std::byte> bs = as_bytes(s);
+      Span<float, is_device> s;
+      const Span<const std::byte, is_device> bs = as_bytes(s);
       SPAN_ASSERT_TRUE(bs.size() == s.size(), status_);
       SPAN_ASSERT_TRUE(bs.size() == 0, status_);
       SPAN_ASSERT_TRUE(bs.size_bytes() == 0, status_);
@@ -177,6 +200,7 @@ struct TestAsBytes {
   }
 };
 
+template <bool is_device>
 struct TestAsWritableBytes {
   int* status_;
 
@@ -188,20 +212,21 @@ struct TestAsWritableBytes {
     InitializeRange(arr, arr + 16);
 
     {
-      Span<float> s;
-      Span<std::byte> bs = as_writable_bytes(s);
-      SPAN_ASSERT_TRUE(bs.size() == s.size(), status_);
-      SPAN_ASSERT_TRUE(bs.size_bytes() == s.size_bytes(), status_);
-      SPAN_ASSERT_TRUE(bs.size() == 0, status_);
-      SPAN_ASSERT_TRUE(bs.size_bytes() == 0, status_);
-      SPAN_ASSERT_TRUE(bs.data() == nullptr, status_);
+      Span<float, is_device> s;
+      Span<std::byte, is_device> byte_s = as_writable_bytes(s);
+      SPAN_ASSERT_TRUE(byte_s.size() == s.size(), status_);
+      SPAN_ASSERT_TRUE(byte_s.size_bytes() == s.size_bytes(), status_);
+      SPAN_ASSERT_TRUE(byte_s.size() == 0, status_);
+      SPAN_ASSERT_TRUE(byte_s.size_bytes() == 0, status_);
+      SPAN_ASSERT_TRUE(byte_s.data() == nullptr, status_);
       SPAN_ASSERT_TRUE(
-        static_cast<void*>(bs.data()) == static_cast<void*>(s.data()), status_);
+        static_cast<void*>(byte_s.data()) == static_cast<void*>(s.data()),
+        status_);
     }
 
     {
-      Span<float> s{arr};
-      Span<std::byte> bs{as_writable_bytes(s)};
+      Span<float, is_device> s{arr};
+      Span<std::byte, is_device> bs{as_writable_bytes(s)};
       SPAN_ASSERT_TRUE(s.size_bytes() == bs.size_bytes(), status_);
       SPAN_ASSERT_TRUE(
         static_cast<void*>(bs.data()) == static_cast<void*>(s.data()), status_);