 #include <thrust/copy.h>
 #include <thrust/sort.h>
-#include <thrust/device_vector.h>
+#include <thrust/memory.h>
+#include <thrust/system/cuda/memory.h>
 
 #include <new> // for std::bad_alloc
 #include <cassert>

  {
    for (size_t n = (1 << 20); n < ((size_t) 1 << 40); n *= 2)
    {
      // TODO ideally we'd use the fallback_allocator in the device_vector too
      //thrust::device_vector<int, fallback_allocator> d_vec(n);

      std::cout << "attempting to sort " << n << " values" << std::endl;


      {
        kernel<<<100,256>>>(raw_ptr, n); // generate unsorted values

        thrust::pointer<int,my_tag> begin = thrust::pointer<int,my_tag>(raw_ptr);
        thrust::pointer<int,my_tag> end   = begin + n;

        thrust::sort(begin, end);

        assert(thrust::is_sorted(begin, end));

