             pq->d_queue[selector][my_valid+far_pile_offset] = my_vert;
         
     }
+
+    static __device__ __forceinline__ void Compact2(
+                                            VertexId        *&vertex_in,
+                                            NearFarPile     *&pq,
+                                            int             &selector,
+                                            SizeT           &input_queue_length,
+                                            VertexId        *&vertex_out,
+                                            SizeT           &v_out_offset,
+                                            SizeT           &node_num)
+    {
+        int tid = threadIdx.x;
+        int bid = blockIdx.x;
+        int my_id = bid*blockDim.x + tid;
+        if (my_id >= input_queue_length)
+            return;
+
+        unsigned int my_vert = vertex_in[my_id];
+        if (my_vert<0 || my_vert >= node_num) return;
+        unsigned int my_valid = pq->d_valid_near[my_id];
+
+        if (my_valid == pq->d_valid_near[my_id+1]-1)
+            vertex_out[my_valid+v_out_offset] = my_vert;
+    }
 };
 
 /**

        //printf("valid near, far: %d, %d\n", pq->d_valid_near[my_id], pq->d_valid_far[my_id]);
    }

    static __device__ __forceinline__ void Compact(
                                            VertexId        *&vertex_in,
                                            NearFarPile     *&pq,

    return close_size[0];
}

/**
 * @brief Remove redundant elements in a queue
 *
 * @tparam KernelPolicy Kernel policy type.
 * @tparam ProblemData Problem data type.
 * @tparam PriorityQueue PriorityQueue data type.
 * @tparam Functor Functor type.
 *
 * @param[in] vertex_in     Device pointer of the input vertex IDs
 * @param[out] pq           PriorityQueue pointer which will be used to store the near/far pile after the input vertices is splitted.
 * @param[in] input_queue_length Input queue length
 * @param[in] problem       Problem object which stores user-defined priority value
 * @param[out] vertex_out   Device pointer of the output vertex IDs, will be used for any other following operators
 * @param[in] far_pile_offset Where to append the newly generated elements in the far pile
 * @param[in] lower_limit   Near pile priority value threshold
 * @param[in] upper_limit   Far pile priority value threshold
 * @param[in] context       CudaContext pointer for moderngpu APIs
 */
template <typename KernelPolicy, typename ProblemData, typename PriorityQueue, typename Functor>
    unsigned int RemoveInvalid(
        typename KernelPolicy::VertexId     *vertex_in,
        PriorityQueue                       *pq,
        typename KernelPolicy::SizeT        input_queue_length,
        typename ProblemData::DataSlice     *problem,
        typename KernelPolicy::VertexId     *vertex_out,
        CudaContext                         &context,
        typename KernelPolicy::SizeT        node_num)
{

    typedef typename KernelPolicy::VertexId     VertexId;
    typedef typename KernelPolicy::SizeT        SizeT;
    typename PriorityQueue::NearFarPile         *nf_pile = pq->d_nf_pile[0];

    int block_num = (input_queue_length + KernelPolicy::THREADS - 1) / KernelPolicy::THREADS;
    unsigned int close_size[1];
    unsigned int far_size[1];
    close_size[0] = 0;
    far_size[0] = 0;
    if(input_queue_length > 0)
    {
        MarkVisit<KernelPolicy, ProblemData, PriorityQueue, Functor><<<block_num, KernelPolicy::THREADS>>>(vertex_in, problem, input_queue_length, node_num);
        // MarkValid
        MarkValid<KernelPolicy, ProblemData, PriorityQueue, Functor><<<block_num, KernelPolicy::THREADS>>>(vertex_in, nf_pile, problem, input_queue_length, node_num);

        // Scan(near)
        // Scan(far)
        Scan<mgpu::MgpuScanTypeExc>(pq->nf_pile[0]->d_valid_near, input_queue_length+1, 0, mgpu::plus<VertexId>(), (VertexId*)0, (VertexId*)0, pq->nf_pile[0]->d_valid_near, context);
        // Compact2
        Compact2<KernelPolicy, ProblemData, PriorityQueue, Functor><<<block_num, KernelPolicy::THREADS>>>(vertex_in, nf_pile, pq->selector, input_queue_length, vertex_out, 0, node_num);
        // get output_length
        cudaMemcpy(&close_size[0], pq->nf_pile[0]->d_valid_near+input_queue_length, sizeof(VertexId),
		    cudaMemcpyDeviceToHost);
    }
    // Update queue length
    return close_size[0];
}

} //priority_queue
} //gunrock

