         return retval;
     }
 
-    // Init function takes a CSR graph stored on CPU, and initialize graph
-    // topology data (by loading ProblemBase::Init(graph)) and DataSlice on
-    // GPU.
+    // The Init function initializes this Problem struct with a CSR
+    // graph that's stored on the CPU. It also initializes the
+    // algorithm-specific data, here depth and predecessor.
     cudaError_t Init(
             const Csr<VertexId, Value, SizeT> &graph)
     {

        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;  
    }

    // Reset function will be loaded before each BFS to reset problem related data.
    cudaError_t Reset(
            const Csr<VertexId, Value, SizeT> &graph, VertexId src)
    {
        cudaError_t retval = cudaSuccess;
        if (retval = util::GRError(ProblemBase::Reset(graph))) break;
        // Reset depth values and predecessor values to invalid. Only set 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 source node ID into the initial frontier
        if (retval = util::GRError(CopyGPU2CPU(g_slices[0]->ping_pong_working_queue, src, 1)));
        return retval;  
    }
};

    SizeT edges;
    DataSlice *d_data_slices;

    // The constructor and destructor are ignored here.

    // "Extract" copies labels and predecessors back to the CPU.
    cudaError_t Extract(VertexId *h_labels, VertexId *h_preds)
    {
        cudaError_t retval = cudaSuccess;
