             d_neighbors, d_adj, nneigh,
             thrust::raw_pointer_cast(out_neighbor_ptr),
             *mgpu_context);
+
+        // sort tags by neighbors
+        mgpu::LocalitySortPairs(thrust::raw_pointer_cast(out_neighbor_ptr),
+                       d_ghost_tag, n_out, *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(plan_out_ptr),0);
        }

    thrust::lower_bound(thrust::cuda::par(alloc),
        out_neighbor_ptr,
        out_neighbor_ptr + n_out,

    thrust::device_ptr<unsigned int> out_neighbor_ptr((unsigned int *)alloc.allocate(n_out*sizeof(unsigned int)));

    // Optimization: are we sending into more than two independent directions simultaneously?
    if (true || num > 1)
        {
        /*
         * expand each tag by the number of neighbors to send the corresponding ptl to
