-// Problem struct is inherited from ProblemBase struct, which stores graph
-// topology data in CSR format and frontier buffers.
+// This data structure (the "Problem" struct) stores the graph
+// topology in CSR format and the frontier. All Problem structs
+// inherit from the ProblemBase struct. Algorithm-specific data is
+// stored in a "DataSlice".
 
 template<
          typename VertexId,

         struct BFSProblem : public ProblemBase<VertexId, SizeT, _USE_DOUBLE_BUFFER>
{

    // MARK_PREDECESSOR would be true when algorithm needs to know the predecessor node IDs
    // for a newly formed frontier during a certain step.
    static const bool MARK_PREDECESSORS     = _MARK_PREDECESSORS;
    // ENABLE_IDEMPOTENCE would be true when the operation performed in parallel
    // for all neighbor nodes/edges is idempotent, meaning data race is benign.
    static const bool ENABLE_IDEMPOTENCE    = _ENABLE_IDEMPOTENCE;

    // DataSlice sturct stores per-node or per-edge arrays and global variables (if any).
    struct DataSlice
    {
        VertexId *d_labels; // BFS depth value

        if (retval = util::GRError(ProblemBase::Init(graph))) break;
        if (retval = util::GRError(GPUMalloc(data_slices[0]->d_labels, nodes))) break;
        if (retval = util::GRError(GPUMalloc(data_slices[0]->d_preds, nodes))) break;
        return retval;
    }

    // The Reset function primes the graph data structure to an
    // untraversed state.
    cudaError_t Reset(
            const Csr<VertexId, Value, SizeT> &graph, VertexId src)
    {
        cudaError_t retval = cudaSuccess;
        if (retval = util::GRError(ProblemBase::Reset(graph))) break;
        // Set all depth and predecessor values to invalid. Set the
        // source node's depth value to 0.
        util::MemsetKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_labels, INVALID_NODE_VALUE, nodes);
        util::MemsetKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_preds, INVALID_PREDECESSOR_ID, nodes);
        if (retval = util::GRError(CopyGPU2CPU(data_slices[0]->d_labels+src, 0, 1)));
        // Put the source node ID into the initial frontier.
        if (retval = util::GRError(CopyGPU2CPU(g_slices[0]->ping_pong_working_queue, src, 1)));
        return retval;
    }
};
