       {
         kernel<<<100,256>>>(raw_ptr, n); // generate unsorted values
 
-        thrust::device_ptr<int> begin = thrust::device_pointer_cast(raw_ptr);
-        thrust::device_ptr<int> end   = begin + n;
+        thrust::pointer<int,my_tag> begin = thrust::pointer<int,my_tag>(raw_ptr);
+        thrust::pointer<int,my_tag> end   = begin + n;
 
-        thrust::sort(thrust::retag<my_tag>(begin), thrust::retag<my_tag>(end));
+        thrust::sort(begin, end);
 
         assert(thrust::is_sorted(begin, end));
 

#include <thrust/copy.h>
#include <thrust/sort.h>
#include <thrust/device_vector.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 vector too
      //thrust::cuda::vector<int, fallback_allocator> d_vec(n);

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

