 #include <thrust/iterator/zip_iterator.h>
 
 // moderngpu
+#include "moderngpu/util/mgpucontext.h"
 #include "moderngpu/device/loadstore.cuh"
 #include "moderngpu/device/launchbox.cuh"
 #include "moderngpu/device/ctaloadbalance.cuh"
-#include "moderngpu/kernels/intervalmove.cuh"
-#include "moderngpu/util/mgpucontext.h"
+#include "moderngpu/kernels/localitysort.cuh"
+#include "moderngpu/kernels/search.cuh"
+#include "moderngpu/kernels/sortedsearch.cuh"
 
 using namespace thrust;
 

            d_neighbors, d_adj, nneigh,
            thrust::raw_pointer_cast(out_neighbor_ptr),
            *mgpu_context);
        #endif

        alloc.deallocate((char *)thrust::raw_pointer_cast(offsets_ptr),0);
        /*
         * sort by neighbor and compute start and end indices
         */
        thrust::sort_by_key(thrust::cuda::par(alloc),
            out_neighbor_ptr,
            out_neighbor_ptr + n_out,
            ghost_tag_ptr);

        }
    else
        {


        alloc.deallocate((char *)thrust::raw_pointer_cast(out_idx_ptr),0);
        alloc.deallocate((char *)thrust::raw_pointer_cast(n_ptr),0);

        /*
         * sort by neighbor and compute start and end indices
         */
        thrust::sort_by_key(thrust::cuda::par(alloc),
            out_neighbor_ptr,
            out_neighbor_ptr + n_out,
            ghost_tag_ptr);


        #else
        gpu_expand_neighbors(n_out,
            thrust::raw_pointer_cast(offsets_ptr),
