bfs_enactor.cuh

#

The enactor defines how a graph primitive runs. It calls traversal (advance and filter operators) and computation (functors).

#
class BFSEnactor : public EnactorBase {
#

For BFS, Constructor, Destructor, and Setup functions are ignored

    template<
    typename AdvancePolicy,
    typename FilterPolicy,
    typename BFSProblem>
    cudaError_t EnactBFS(CudaContext &context,BFSProblem *problem,VertexId   src)
    {

        typedef BFSFunctor<
        typename BFSProblem::VertexId,
        typename BFSProblem::SizeT,
        typename BFSProblem::VertexId,
        BFSProblem> BfsFunctor;
#

Start with the Setup function to initialize kernel running parameters

        cudaError_t retval = cudaSuccess;
        if (retval = EnactorBase::Setup(problem)) break;
#

Define the graph topology data pointer (g_slice) and the problem-specific data pointer (d_slice)

        typename BFSProblem::GraphSlice *g_slice = problem->d_graph_slices;
        typename BFSProblem::DataSlice *d_slice = problem->d_data_slices;
#

Initialize the queue length (frontier size) to 1.

        SizeT queue_length = 1;
#

We ping-pong between old and new frontiers; "selector" picks which one is the current destination.

        int selector = 0;
#

Here we sequence our operators and functors. For BFS, we alternate between advancing to a new frontier (vertex-to-vertex), calling the BFS functor to set depths along the way, and then filtering out invalid nodes from the new frontier. We repeat until the frontier is empty.

        while (queue_length > 0) {
            gunrock::oprtr::advance::Kernel
                <AdvancePolicy, BFSProblem, BFSFunctor>
                <<<advance_grid_size, AdvancePolicy::THREADS>>>(
                        queue_length,
                        g_slice->ping_pong_working_queue[selector],
                        g_slice->ping_pong_working_queue[selector^1],
                        d_slice,
                        context,
#

This advance is vertex-to-vertex

                        gunrock::oprtr::advance::V2V);

            selector ^= 1; // Swap selector

            gunrock::oprtr::filter::Kernel
                <FilterPolicy, BFSProblem, BFSFunctor>
                <<<filter_grid_size, FilterPolicy::THREADS>>>(
                        queue_length,
                        g_slice->ping_pong_working_queue[selector],
                        g_slice->ping_pong_working_queue[selector^1],
                        d_slice);
        }
        return retval;
    }
#

The entry point in the driver code to BFS is this Enact call.

    template <typename BFSProblem>
        cudaError_t Enact(
                CudaContext &context,
                BFSProblem *problem, // Problem data sent in
                typename BFSProblem::VertexId src) // Source node ID for BFS
        {
#

Gunrock provides recommended settings here for kernel parameters, but they can be changed by end-users.

            typedef gunrock::oprtr::filter::KernelPolicy<
                BFSProblem,
                300, //CUDA_ARCH
                8, //MIN_CTA_OCCUPANCY
                8> //LOG_THREAD_NUM
                FilterKernelPolicy;

            typedef gunrock::oprtr::advance::KernelPolicy<
                BFSProblem,
                300,//CUDA_ARCH
                8, //MIN_CTA_OCCUPANCY
                10, //LOG_THREAD_NUM
                32*128> //THRESHOLD_TO_SWITCH_ADVANCE_MODE
                AdvanceKernelPolicy;

            return EnactBFS<AdvanceKernelPolicy, FilterKernelPolicy, BFSProblem>(context, problem, src);
        }
};