     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(), 

                          valuesBegin, cubeOwners.begin(), Begin(), 
                          thrust::equal_to<int>(), ReduceHyperCubes());
    CHECK_FOR_CUDA_ERROR();
    
#else
    // Reduce on the CPU
    thrust::host_vector<unsigned int> hostRayOwners(rayOwners.begin(), rayOwners.end());
    thrust::host_vector<SignedAxis> hostA = A;
    thrust::host_vector<float2> hostX = X;
    thrust::host_vector<float2> hostY = Y;
    thrust::host_vector<float2> hostZ = Z;
    thrust::host_vector<float2> hostU = U;
    thrust::host_vector<float2> hostV = V;
    
    thrust::zip_iterator<thrust::tuple<thrust::host_vector<SignedAxis>::iterator,
        thrust::host_vector<float2>::iterator,
        thrust::host_vector<float2>::iterator,
        thrust::host_vector<float2>::iterator,
        thrust::host_vector<float2>::iterator,
        thrust::host_vector<float2>::iterator> > hostValuesBegin = 
        thrust::make_zip_iterator(thrust::make_tuple(hostA.begin(), hostX.begin(), hostY.begin(), 
                                                     hostZ.begin(), hostU.begin(), hostV.begin()));
    
    static thrust::host_vector<int> cubeOwners(128); cubeOwners.resize(Size());

    thrust::host_vector<SignedAxis> hostResA(Size());
    thrust::host_vector<float2> hostResX(Size());
    thrust::host_vector<float2> hostResY(Size());
    thrust::host_vector<float2> hostResZ(Size());
    thrust::host_vector<float2> hostResU(Size());
    thrust::host_vector<float2> hostResV(Size());
    
    thrust::zip_iterator<thrust::tuple<thrust::host_vector<SignedAxis>::iterator,
        thrust::host_vector<float2>::iterator,
        thrust::host_vector<float2>::iterator,
        thrust::host_vector<float2>::iterator,
        thrust::host_vector<float2>::iterator,
        thrust::host_vector<float2>::iterator> > hostCubeVals = 
        thrust::make_zip_iterator(thrust::make_tuple(hostResA.begin(), 
                                                     hostResX.begin(), 
                                                     hostResY.begin(), 
                                                     hostResZ.begin(), 
                                                     hostResU.begin(), 
                                                     hostResV.begin()));
    
    thrust::reduce_by_key(hostRayOwners.begin(), hostRayOwners.end(),
                          hostValuesBegin, cubeOwners.begin(), hostCubeVals, 
                          thrust::equal_to<int>(), ReduceHyperCubes());
    
    a = hostResA;
    x = hostResX;
    y = hostResY;
    z = hostResZ;
    u = hostResU;
    v = hostResV;
#endif
}

size_t HyperCubes::Resize(const size_t s) {

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

#include <HyperCube.h>

