 // into the device address space.  The fallback_allocator enables
 // the GPU to process data sets that are larger than the device
 // memory, albeit with a significantly reduced performance.
-
-
-// initialize some unsorted data
-__global__
-void kernel(int * d_ptr, size_t N)
-{
-  size_t thread_id = blockDim.x * blockIdx.x + threadIdx.x;
-  size_t grid_size = blockDim.x * gridDim.x;
-
-  for (size_t i = thread_id; i < N; i += grid_size)
-    d_ptr[i] = i % 1024;
-}
 
 
 // derive a simple allocator from cuda::dispatchable for using pinned host memory as a functional fallback


      if(raw_ptr)
      {
        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;

        // sort the data using our special allocator
        // if temporary memory is required during the sort,

      {
        // attempt to deallocate buffer
        std::cout << "  failed to map host memory into device address space (fallback failed)" << std::endl;
        cudaError_t error = cudaFreeHost(h_ptr);
        if(error)
        {
          throw thrust::system_error(error, thrust::cuda_category(), "Error after cudaFreeHost");
        }

        result = 0;
      }
    }
