         i != free_blocks.end();
         ++i)
     {
-      // transform the pointer to device_ptr before calling device_free
-      thrust::device_free(thrust::device_pointer_cast(i->second));
+      // transform the pointer to cuda::pointer before calling cuda::free
+      thrust::system::cuda::free(thrust::system::cuda::pointer<void>(i->second));
     }
 
     for(allocated_blocks_type::iterator i = allocated_blocks.begin();
         i != allocated_blocks.end();
         ++i)
     {
-      // transform the pointer to device_ptr before calling device_free
-      thrust::device_free(thrust::device_pointer_cast(i->first));
+      // transform the pointer to cuda::pointer before calling cuda::free
+      thrust::system::cuda::free(thrust::system::cuda::pointer<void>(i->first));
     }
   }
 

  // generate random input
  thrust::generate(h_input.begin(), h_input.end(), rand);

  thrust::device_vector<int> d_input = h_input;
  thrust::device_vector<int> d_result(n);

  size_t num_trials = 5;


// during algorithms such as thrust::sort. The idea will be to create a simple
// cache of allocations to search when temporary storage is requested. If a hit
// is found in the cache, we quickly return the cached allocation instead of
// resorting to the more expensive thrust::system::cuda::malloc.
//
// Note: this implementation cached_allocator is not thread-safe. If multiple
// (host) threads use the same cached_allocator then they should gain exclusive
