   cudaGetDevice(&device);
   cudaDeviceProp properties;
   cudaGetDeviceProperties(&properties, device);
+
+  fallback_allocator alloc;
 
   if (!properties.canMapHostMemory)
   {

};


// the fallback allocator is simply a global variable
// XXX ideally this variable is declared thread_local
fallback_allocator g_allocator;

// create a tag derived from cuda::tag for distinguishing
// our overloads of get_temporary_buffer and return_temporary_buffer
struct my_tag : thrust::cuda::tag {};


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

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

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



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

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

        thrust::sort(alloc, begin, end);

        robust_cudaFree(raw_ptr);
      }
