salsa_problem.cuh

#

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,
typename SizeT,
typename Value>
struct SALSAProblem : public ProblemBase
{

    static const bool MARK_PREDECESSORS     = true; // for SALSA algorithm, we need to track predecessors in Advance
    static const bool ENABLE_IDEMPOTENCE    = false; // In SALSA, data races during Advance are not allowed.

    struct DataSlice
    {
        Value *d_hrank_curr; // hub rank score for current iteration
        Value *d_arank_curr; // authority rank score for current iteration
        Value *d_hrank_next; // hub rank score for next iteration
        Value *d_arank_next; // authority rank score for next iteration
        VertexId *d_in_degrees; // in degrees for each node
        VertexId *d_out_degrees; // out degrees for each node
        VertexId *d_hub_predecessors; // hub graph predecessors (original graph)
        VertexId *d_auth_predecessors; // authority graph predecessors (reverse graph)
    };

    SizeT nodes; // node number of the graph
    SizeT edges; // edge number of the graph
    SizeT out_nodes; // number of nodes that have outgoing edges
    SizeT in_nodes; // number of nodes that have incoming edges
    DataSlice *d_data_slices;
#

The constructor and destructor are ignored here.

#

"Extract" copies final hub rank scores and authority rank scores back to the CPU.

    cudaError_t Extract(VertexId *h_hrank, VertexId *h_arank)
    {
        cudaError_t retval = cudaSuccess;
        if (retval = util::GRError(CopyGPU2CPU(data_slices[0]->d_hrank_curr, h_hrank, nodes))) break;
        if (retval = util::GRError(CopyGPU2CPU(data_slices[0]->d_arank_curr, h_arank, nodes))) break;
        return retval;
    }
#

SALSA is an algorithm for bipartite graphs. So the Init function takes two CSR graphs---the hub (original) graph and the auth (reverse) graph stored on the CPU, and uses them to initialize the graph topology data and the SALSA-specific DataSlice on the GPU.

    cudaError_t Init(
            const Csr<VertexId, Value, SizeT> &hub_graph,
            const Csr<VertexId, Value, SizeT> &auth_graph)
    {
        cudaError_t retval = cudaSuccess;
        if (retval = util::GRError(ProblemBase::Init(hub_graph, auth_graph))) break;
        if (retval = util::GRError(GPUMalloc(data_slices[0]->d_hrank_curr, nodes))) break;
        if (retval = util::GRError(GPUMalloc(data_slices[0]->d_arank_curr, nodes))) break;
        if (retval = util::GRError(GPUMalloc(data_slices[0]->d_hrank_next, nodes))) break;
        if (retval = util::GRError(GPUMalloc(data_slices[0]->d_arank_next, nodes))) break;
        if (retval = util::GRError(GPUMalloc(data_slices[0]->d_in_degrees, nodes))) break;
        if (retval = util::GRError(GPUMalloc(data_slices[0]->d_out_degrees, nodes))) break;
        if (retval = util::GRError(GPUMalloc(data_slices[0]->d_hub_predecessors, edges))) break;
        if (retval = util::GRError(GPUMalloc(data_slices[0]->d_auth_predecessors, edges))) break;
        return retval;
    }
#

The Reset function primes the graph data structure to an initial state, which includes ...

    cudaError_t Reset(const Csr<VertexId, Value, SizeT> &graph)
    {
        cudaError_t retval = cudaSuccess;
        if (retval = util::GRError(ProblemBase::Reset(graph))) break;
#

... initializing the hub and authority rank scores ...

        util::MemsetKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_hrank_curr, (Value)1.0/out_nodes, nodes);
        util::MemsetKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_arank_curr, (Value)1.0/in_nodes, nodes);
        util::MemsetKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_hrank_next, 0, nodes);
        util::MemsetKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_arank_next, 0, nodes);
#

... accurate in and out degrees for each node ...

        util::MemsetKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_out_degrees, 0, nodes);
        util::MemsetKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_in_degrees, 0, nodes);
        util::MemsetMadVectorKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_out_degrees, BaseProblem::graph_slices[gpu]->d_row_offsets, &BaseProblem::graph_slices[gpu]->d_row_offsets[1], -1, nodes);
        util::MemsetMadVectorKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_in_degrees, BaseProblem::graph_slices[gpu]->d_column_offsets, &BaseProblem::graph_slices[gpu]->d_column_offsets[1], -1, nodes);
#

... and initializing predecessors with an INVALID_PREDECESSOR_ID.

        util::MemsetKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_hub_predecessors, INVALID_PREDECESSOR_ID, edges);
        util::MemsetKernel<<<BLOCK, THREAD>>>(data_slices[0]->d_auth_predecessors, INVALID_PREDECESSOR_ID, edges);

        return retval;
    }
};