 
       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;
+
+        // generate unsorted values
+        thrust::tabulate(begin, end, thrust::placeholders::_1 % 1024);
 
         // 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;
        cudaFreeHost(h_ptr);
        result = 0;
      }
    }

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


// derive a simple allocator from cuda::dispatchable for using pinned host memory as a functional fallback
