 // 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

#include <thrust/copy.h>
#include <thrust/sort.h>
#include <thrust/memory.h>
#include <thrust/system/cuda/memory.h>

#include <new> // for std::bad_alloc
#include <cassert>
#include <iostream>
#include <iterator>

// This example demonstrates how to intercept calls to malloc
// and free to implement a fallback for cudaMalloc.

      {
        // 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;
      }
    }
