     PrepReduceData prepReduceData;
     thrust::transform(rayBegin, rayEnd, transBegin, prepReduceData);
 
-    // TODO HACK Move everything to the host because reduce_by_key fails on 330M
     // TODO Remove owners, at the least reduce by a partition start flag instead
     // e.g. |1|0|0|1|0|0|0|
     static thrust::device_vector<unsigned int> rayOwners(rayRange); rayOwners.resize(rayRange);
     DacrtNodes::CalcOwners(rayPartitions.begin(), rayPartitions.end(), rayOwners);
 
-#if REDUCE_ON_DEVICE
     // Reduce on the GPU
     Iterator valuesBegin = 
         thrust::make_zip_iterator(thrust::make_tuple(A.begin(), X.begin(), Y.begin(), 

#include <iomanip>

#include <thrust/reduce.h>

#if REDUCE_ON_DEVICE<1
#include <thrust/host_vector.h>
#endif

// *** HyperCube ***


// This program is open source and distributed under the New BSD License. See
// license for more detail.
// -----------------------------------------------------------------------------

#include <HyperCube.h>

