       {
         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::cuda::pointer<int> begin = thrust::cuda::pointer<int>(raw_ptr);
+        thrust::cuda::pointer<int> end   = begin + n;
 
-        thrust::sort(begin, end);
+        thrust::sort(alloc, begin, end);
 
         robust_cudaFree(raw_ptr);
       }


// build a simple allocator for using pinned host memory as a functional fallback
struct fallback_allocator
{
  fallback_allocator() {}


};


// overload get_temporary_buffer on fallback_allocator
// its job is to forward allocation requests to alloc
template<typename T>
  thrust::pair<T*, std::ptrdiff_t>
    get_temporary_buffer(fallback_allocator &alloc, std::ptrdiff_t n)
{
  // ask the allocator for sizeof(T) * n bytes
  T* result = reinterpret_cast<T*>(alloc.allocate(sizeof(T) * n));

  // return the pointer and the number of elements allocated
  return thrust::make_pair(result,n);
}

// overload return_temporary_buffer on fallback_allocator
// its job is to forward deallocations to alloc
// an overloaded return_temporary_buffer should always accompany
// an overloaded get_temporary_buffer
template<typename Pointer>
  void return_temporary_buffer(fallback_allocator &alloc, Pointer p)
{
  // return the pointer to the allocator
  alloc.deallocate(thrust::raw_pointer_cast(p));
}


