GraphBLAST: A High-Performance Linear Algebra-based Graph Framework on the GPU

2.1 Related Work

Large-scale graph frameworks on multi-threaded CPUs, distributed-memory CPU systems (surveyed by Batarfi et al. [3]), and massively parallel GPUs (surveyed by Shi et al. [74]) fall into three broad categories: vertex-centric (surveyed by McCune, Weninger, and Madey [58]), edge-centric, and linear-algebra-based. In this section, we will explain this categorization and the influential graph frameworks from each category.

2.2 Previous Systems

Two systems that directly inspired our work are Gunrock and Ligra.

2.3 Graph Traversal vs. Matrix-vector Multiply

The connection between graph traversal and linear algebra was noted by Denes König [54] in the early days of graph theory. Since then the connection between graphs and matrices has been established by the popular representation of a graph as an adjacency matrix. More specifically, it has become popular to represent a vector-matrix multiply as being equivalent to one iteration of breadth-first-search traversal (see Figure 2). Some seemingly non-traversal graph algorithms such as triangle counting, which can be solved efficiently using a masked SpGEMM, can also be thought in terms of traversals. Multiplying the lower triangle of the adjacency matrix with its transpose, as we do in Section 7.5, simply does a 2-hop traversal of the graph from every vertex. The use of lower triangle and its transpose as opposed to the whole adjacency matrix ensures that identical paths are not explored redundantly. The masking checks whether such 2-hop paths, also called wedges or triads in the literature, close to form triangles.

Fig. 2.

View Figure

3.1 Abstract Algebraic Constructs

3.2 Programming Constructs

3.3 Operation

An operation is a commonly used linear algebra operation. A full list of operations is shown in Table 7. Of the operations in Table 7, the most computationally intensive and useful operations are mxm, mxv, and vxm. We find empirically that these operations take over 90% of graph algorithm runtime. For these operations, we decompose them into constituent parts in order to better optimize their performance (see Figure 4). Intuitively, mxv and vxm represent a graph traversal from a single-nodeset (one single set of active nodes \(\mathbf {u}(i)\) for all i). On ther other hand, mxm represents a multi-nodeset, that is to say, a graph-traversal from multiple, independent sets of active nodes at the same time (for all independent sets of active nodes \(B(:,j)\), start a traversal for all active nodes i in that set).

An operation is used to tie all the objects in GraphBLAS together. Figure 3 shows a matrix-vector multiply, which represents a graph traversal from all active nodes \(\mathbf {u}(i)\) to their neighbor nodes, applying the semiring multiply \(\otimes\) to get a temporary \(\mathbf {c}(i,j) = \mathbf {A}(i, j) \otimes \mathbf {u}(i)\) between the edge value \(\mathbf {A}(i,j)\) and the active node’s value \(\mathbf {u}(i)\) and finally doing a reduction over the temporary output \(\mathbf {w}(i) = \bigoplus _j \mathbf {c}(i,j)\) using the semiring add operation \(\oplus\) and the semiring identity \(\mathrm{I}\). Optionally, the user has the option of applying a Boolean mask to the output so that the output is effectively \(\mathbf {w}(i) = \mathbf {mask}(i) \times \mathbf {w}(i)\) for all i. The descriptor can further mutate this operation to account for matrix transposition and whether the mask or its structural complement should be used.

Fig. 3.

View Figure

3.4 Running Example

As a running example in this paper, we discuss sparse-matrix multiplication by dense-vector (SpMV) and sparse-vector (SpMSpV) with a direct dependence on graph traversal. The key step we will be discussing is Line 8 of Algorithm 1, which is the matrix formulation of parallel breadth-first-search. Illustrated in Figure 5(b), this problem consists of a matrix-vector multiply followed by an elementwise multiplication between two vectors: one is the output of the matrix-vector multiply and the other is the negation (or structural complement) of the visited vector.

Using the standard dense matrix-vector multiplication algorithm (GEMV), we would require \(8\ \times \ 8 = 64\) loads and store instructions. However, if we instead treat the matrix not as a dense matrix but as a sparse matrix in order to take advantage of input matrix sparsity, we can perform the same computation in a mere 20 loads and stores into the sparse matrix. This number comes from counting the number of nonzeroes in the sparse matrix, which is equivalent to the number of edges in the graph. Using this as the baseline, we will show in later sections how optimizations such as exploiting the input vector and output vector sparsity can further reduce the number of loads and stores required.

3.5 Code Example

Having described the different components of GraphBLAST, we show a short code example of how to do breadth-first-search using the GraphBLAST interface alongside the linear algebra in Algorithm 1. Before the while-loop, the vectors \(\mathbf {f}\) and \(\mathbf {v}\) (representing the vertices currently active in the traversal and the set of previously visited vertices) are initialized.

Then in each iteration of the while-loop, the following steps take place: (1) vertices currently active are added to the visited vertex vector, marked by the iteration d where they were first encountered; (2) the active vertices are traversed to find the next set of active vertices, and then elementwise-multiplied by the negation of the set of active vertices (filtering out previously visited vertices); (3) the number of active vertices of the next iterations is reduced to variable c; (4) the iteration number is incremented. This while-loop continues until there are no more active vertices (c reaches 0).

As demonstrated in the code example, GraphBLAS has the advantage of being concise. Developing new graph algorithms in GraphBLAS requires modifying a single file and writing simple C++ code. Provided a GraphBLAS implementation exists for a particular processor, GraphBLAS code can be used with minimal changes. Over the next three sections, we will discuss the most important design principles for making this code performant: exploiting input sparsity, exploiting output sparsity, and good load-balancing.

3.6 Differences with GraphBLAS C API Standard

We have tried to make our framework interface adhere to the GraphBLAS C API as close as possible, but since we decided to take advantage of certain C++ features (templates and functors, in particular) in our framework, we have departed from the strict API standard. We observe that some differences are motivation for the design of a GraphBLAS C++ API [14]. Some major differences in the interface are listed below:

A few minor interface differences follow:

4.1 Two Roads to Matrix-vector Multiplication

Before we address the above challenges, we draw a distinction between two different ways the matrix-vector multiply \(\mathbf {y} \leftarrow \mathbf {A} \mathbf {x}\) can be computed. We distinguish between multiplying a sparse-matrix by a dense-vector (SpMV) and by a sparse-vector (SpMSpV). There is extensive literature focusing on SpMV for GPUs (including a comprehensive survey [35]). However, we concentrate on SpMSpV, because it is more relevant to graph search algorithms where the vector represents the subset of vertices that are currently active and is typically sparse.

Recall in the running example in Section 3.4 that by exploiting matrix sparsity (SpMV) in favor of dense matrix-vector multiplication (GEMV), we were able to bring the number of load and store instructions down from GEMV’s 64 to SpMV’s 20. A natural question to ask is whether it is possible to decrease the number of load and store instructions further when the input vector is sparse. Indeed, when we exploit input sparsity (SpMSpV) to get the situation in Figure 6(b), we can reduce the number of loads and stores from 20 to 8. Similar to how our move from GEMV to SpMV involved changing the matrix storage format from dense to sparse, moving from SpMV to SpMSpV motivates storing the vector in sparse format. It is worth noting that the sparse vectors are assumed to be implemented as lists of indices and values. A summary is shown in Table 8.

Fig. 6.

View Figure

4.2 Related Work

Mirroring the dichotomy between SpMSpV and SpMV, there are two methods to perform one iteration of graph traversal, which are called push and pull.² They can be used to describe graph traversals in a variety of graph traversal-based algorithms such as breadth-first-search, single-source shortest-path, and PageRank.

In the case of breadth-first-search, push begins with the current frontier (the set of vertices from which we are traversing the graph) and looks for children of this set of vertices. Then, from this set of children, the previously visited vertices must be filtered out to generate the output frontier (the frontier we use as input on the next iteration). In contrast, pull starts from the set of unvisited vertices and looks back to find their parents. If a node in the unvisited-vertex set has a parent in the current frontier, we add it to the output frontier. Beamer, Asanović, and Patterson [7] observed that in the middle iterations of a BFS on scale-free graphs, the frontier becomes large and each neighbor is found many times, leading to redundant work. They show that for optimal performance, in these intermediate iterations, they should switch to pull, and then in later iterations, back to push.

Many graph algorithms such as breadth-first-search, single-source shortest-path, and PageRank involve multiple iterations of graph traversals. Switching between push and pull in different iterations applied to the specific algorithm of breadth-first-search is called direction optimization or direction-optimized BFS, which was also described by Beamer, Asanović, and Patterson [7]. This approach is also termed push-pull. Building on this work, Shun and Blelloch [76] generalized direction optimization to graph traversal algorithms beyond BFS. To avoid confusion with the BFS-specific instance, we refer to Shun and Blelloch’s contribution as generalized direction optimization.

Beamer, Asanović and Patterson later studied matrix-vector multiplication in the context of SpMV- and SpMSpV-based implementations for PageRank [8]. In both their work and that of Besta et al. [11], the authors noted that switching between push/pull is the same as switching between SpMSpV/SpMV. In both works, authors show a one-to-one correspondence between push and SpMSpV, and between pull and SpMV; they are two ways of thinking about the same concept.

Our work differs from Beamer, Asanović, and Patterson and Besta et al. in three ways: (1) they emphasize graph algorithm research, whereas we focus on building a graph framework, (2) their work targets multithreaded CPUs, while ours targets GPUs, and (3) their interface is vertex-centric, but ours is linear-algebra-based.

The work we present here builds on our earlier work and is first to extend the generalized direction optimization technique to linear-algebra-based frameworks based on the GraphBLAS specification. In contrast, previous implementations to the GraphBLAS specification, such as GBTL [88] and SuiteSparse [27], do not support generalized direction optimization and as a consequence, trail state-of-art graph frameworks in performance.

In both implementations, the operation mxv is implemented as a special case of mxm when one of the matrix dimensions is 1 (i.e., is a Vector). The mxm implementation is a variant of Gustavson’s algorithm [41], which takes advantage of both matrix sparsity and input vector sparsity, so it has a similar performance characteristic as SpMSpV. Therefore, it shares SpMSpV’s poor performance when either: (1) there are more elements in the input vector or (2) when there are fewer elements in the mask (representing fewer operations that need to be performed). In other words, neither GBTL and SuiteSparse automatically switch to pull when the input vector becomes large in the middle iterations of graph traversal algorithms like BFS, and perform push throughout the entire BFS. In comparison, our graph framework balances exploiting input vector sparsity (SpMSpV) with the efficiency of SpMV during iterations of high input vector sparsity. This helps us match or exceed the performance of existing graph frameworks (Section 8).

4.3 Implementation

In this subsection, we revisit the three challenges we claimed boil down to different facets of the same challenge: exploiting input sparsity.

• Our backend automatically handles direction optimization when mxv is called, by consulting an empirical cost model and calling either the SpMV or SpMSpV routine that we expect will result in the fewest memory accesses.

• Both routines are necessary for an efficient implementation of mxv in a graph framework.

• For Matrix, store both CSR and CSC, but give users the option to save memory by only storing one of these two representations. The result is a memory-efficient, performance-inefficient solution. For Vector, since both dense vector and sparse vector are required for the two different routines SpMV and SpMSpV, respectively, we give the backend the responsibility to switch between dense and sparse vector representations. We allow the user to specify the initial storage format of the Matrix and Vector objects.

4.4 Direction Optimization Insights

Exploiting input sparsity is a useful and important strategy for high-performance in graph traversals. We believe that the GraphBLAS interface decision where users do not have to specify whether or not they want to exploit input sparsity is a good one; we showed that instead, users must only write code once using the mxv interface and both forms of SpMV and SpMSpV code can be automatically generated for them by GraphBLAST. In the next section, we will show how the number of memory accesses can also be reduced by exploiting output sparsity.

5.1 Motivation and Applications of Masking

Following the brief introduction to masking in Section 5, the reader may wonder why such an operation is necessary. Masking can be thought of in two ways: (i) masking is a way to fuse an element-wise operation with another operation from Table 7; and (ii) masking allows the user to express for which indices they do and do not require a value before the actual computation is performed. We define this as output sparsity. The former means that masking is a way for the user to tell the framework there is an opportunity for kernel fusion, while the latter is an intuitive way to understand why masking can reduce the number of computations in graph algorithms.

There are several graph algorithms where exploiting output sparsity can be used to reduce the number of computations:

5.2 Microbenchmarks

Similar to our earlier microbenchmark (Section 4.3.2), we benchmark how using masked SpMV and SpMSpV variants of graphblas::mxv performed compared with unmasked SpMV and SpMSpV as a function of mask Vector sparsity for a synthetic undirected Kronecker graph with 2M vertices and 182M edges. For more details of the experiment setup, see Section 8.

As our microbenchmark in Figure 7 illustrates, the masked SpMV variant of graphblas::mxv scales with the sparsity of the mask Vector. However, the masked SpMSpV is unchanged from the unmasked SpMSpV. This too matches the theoretical complexity shown in Table 8, which shows that masked SpMV scales with \(O(d\text{ } nnz(\mathbf {m}))\), where \(\mathbf {m}\) is the mask Vector. However in our implementation, masked SpMSpV only performs the elementwise multiply with the mask after the SpMSpV operation, so it is unable to benefit from the mask’s sparsity. After the columns of the sparse matrix are loaded, it may be possible to use binary search into the mask vector to reduce the number of operations required for the multiway merge, but this still does not asymptotically reduce the number of load and store instructions into the sparse matrix as in the masked SpMV case.

Fig. 7.

View Figure

In the running example, recall in Figure 7(a) that standard SpMV, which performs the matrix-vector multiply followed by the elementwise product with the mask, took 20 load and store instructions. However, when we reverse the sequence of operations by first loading the mask, seeing which elements of the mask are nonzero, and then only doing the matrix-vector multiply for those rows that map to a nonzero mask element, we see that the number of loads and stores drops significantly from 20 down to 10.

5.3 Masking Insights

One simple implementation of masking is to perform the matrix multiplication, and then apply the mask to the output. This approach has the benefit of being straightforward and easy to implement. However, we identify two scenarios in which accessing the mask ahead of the matrix multiplication is beneficial:

6.1 Memory Management

GBTL uses the Thrust template library to provide wrappers around GPU memory that automatically handle CPU-to-GPU and GPU-to-CPU communication. We instead decided to manually manage GPU pointers ourselves. This offers graph algorithm developers (i.e., the users) the same seamless experience of not having to concern themselves with whether the data is on the CPU or the GPU.

For the concrete implementation of each object—SparseVector, DenseVector, SparseMatrix, DenseMatrix, etc.—we keep both a canonical GPU copy and a CPU copy that is allowed to go out-of-date. Upon initialization, we maintain the canonical copy on the GPU at all times and use a flag to keep track whether or not the CPU version differs from the canonical GPU copy and is stale. When operations mutating the GPU copy are run, this flag will get set to true indicating the CPU copy is now stale. For operations that interact with the world external to GraphBLAS such as extract and extractTuples, we will copy data back to the CPU depending on whether the flag tells us the CPU copy is stale or not. If it is not stale, we can return the CPU copy directly. If it is stale, we will copy data back to the CPU and reset the flag false indicating the GPU and CPU copy have equal values.

On GPUs, memory allocation time can be a significant fraction of runtime. Since some operations require temporary memory, we keep a memory pool of already-allocated GPU memory. Currently, we associate this memory pool with the Descriptor object, because we do not find we often require more than one Descriptor. In the future, we will consider changing this to use the Factory pattern [52], which is considered standard design for memory pools.

6.2 Operators

One of the biggest challenges of implementing GraphBLAS is solving the problem of supporting the large cross product of possible functionalities. Consider, for example, Semiring. This operator needs to support 11 built-in GraphBLAS types (and any user-defined type), 22 built-in binary operations, and 8 built-in monoid operations. Even without user-defined types, this comes to a total of 1,459 operators for the 3 GraphBLAS methods that take a Semiring, mxv, vxm, and mxm.

In the literature [14, 27, 88], two ways have been proposed to tackle this problem:

The first method is the approach taken by SuiteSparse [27]. This has the advantage of being a shared object library, which can be linked to using frontends written in interpreted languages. We instead adopt the second method, which has the advantage of not needing to maintain code generation tools and macros, which can be challenging.

To express monoids and semirings, we use __host__ and __device__ functors that overload the function call operator (i.e. operator()). We use a macro that constructs structs composed of one and two of these functors for monoids and semirings respectively. The macro also takes an identity-element input.

6.3 Implementation of Key Primitives and Load Balancing

What follows are the implementation details and load balancing techniques behind the four key primitives SpMSpV, SpMV, SpMM and masked SpGEMM. Load balancing attempts to distribute work equally across the GPU’s processors (threads, warps and blocks). To motivate the need for load balance, consider an implementation that assigns each matrix row to a different processor. Because the number of nonzeroes per row may vary greatly, the amount of work per processor may also vary greatly, leading to inefficient execution. We use the following strategies to implement load-balanced kernels:

6.3.2 SpMV.

For SpMV, we use the nonzero-split implementation of the merge-based algorithm provided by ModernGPU [4]. In benchmarking, we find that performance is similar to the merge path algorithm by Merrill and Garland [59]. The difference between the two algorithms is shown in Figure 8 and may be described as follows:

Fig. 8.

View Figure

(1)

Row split [9]: Assigns an equal number of rows to each processor.

(2)

Merge-based: Performs two-phase decomposition—the first kernel divides work evenly amongst CTAs, then the second kernel processes the work. (a) Nonzero-split [4, 25]: Assign an equal number of nonzeros per processor. Then do a 1-D (1-dimensional) binary search on row offsets to determine at which row to start. (b) Merge path [59]: Assign an equal number of {nonzeros and rows} per processor. This is done by doing a 2-D binary search (i.e., on the diagonal line in Figure 8(c)) over row offsets and nonzero indices of matrix \(\mathbf {A}\).

The merge path algorithm has the advantage of doing well in the pathological case of arbitrarily many empty matrix rows, which can cause an arbitrarily large amount of load imbalance for ModernGPU’s nonzero-split. However, we find that this does not happen often in practice. Hence, GraphBLAST currently uses ModernGPU’s SpMV implementation that relies on the segmented-scan primitive. In theory, segmented-scan-based SpMV has depth logarithmic on the number of nonzeros involved [13], and it can be implemented on a GPU in a work-efficient way [72]. While in practice one has to choose a block size on the GPU, Sengupta et al. [71] showed that this does not increase work and depth of the segmented scan implementation asymptotically.

We show the results of a microbenchmark comparing the merge-based algorithm against cuSPARSE’s implementation of SpMV in Figure 9(a), which we suspect relies on the row split algorithm. The experimental setup is described in Section 8. The right side of the x-axis represents load imbalance where long matrix rows are not divided enough, resulting in some computation resources on the GPU remaining idle while others are overburdened. The left size of the x-axis represents load imbalance where too many computational resources are allocated to each row, so some remain idle. From this figure, it is clear that the merge-based algorithm is superior to the row split algorithm at addressing these two types of load imbalance, despite there being two configurations 512 and 2048 for which row split is faster.

Fig. 9.

View Figure

7.1 Breadth-first-search

Given a source vertex \(s \in V\), a BFS is a full exploration of a graph G that produces a spanning tree of the graph, containing all the edges that can be reached from s, and the shortest path from s to each one of them. We define the depth of a vertex as the number of hops it takes to reach this vertex from the root in the spanning tree. The visit proceeds in steps, examining one BFS level at a time. It uses three sets of vertices to keep track of the state of the visit: the frontier contains the vertices that are being explored at the current depth, next has the vertices that can be reached from frontier, and visited has the vertices reached so far. BFS is one of the most fundamental graph algorithms and serves as the basis of several other graph algorithms.

• Hardwired GPU implementationThe best-known BFS implementation of Merrill et al. [60] achieves its high performance through careful load-balancing, avoidance of atomics, and heuristics for avoiding redundant vertex discovery. Its chief operations are expand (to generate a new frontier) and contract (to remove redundant vertices) phases. Enterprise [55], a GPU-based BFS system, introduces a very efficient implementation that combines the benefits of the direction optimization of Beamer, Asanović and Patterson [7], leverages the adaptive load-balancing workload mapping strategy of Merrill et al., and chooses to not synchronize each BFS iteration, which addresses the kernel launch overhead problem (Section 2.1.4).

• GraphBLAST implementationMerrill et al.’s expand and contract maps nicely to GraphBLAST’s mxv operator with a mask using a Boolean semiring. Like Enterprise, we implement efficient load-balancing (Section 6) and direction optimization, which was described in greater detail in Section 4. We do not use Enterprise’s method of skipping synchronization between BFS iterations, but we use two optimizations: early-exit and structure-only, which are consequences of the Boolean semiring that is associated with BFS. We also use operand reuse, which avoids having to convert from sparse to dense during direction optimization. These optimizations are inspired by Gunrock and are described in detail by the authors in an earlier work [85]. If our implementation were to always choose the top-down direction, it would have work complexity of \(O({\it nnz}(\mathbf {A})) = O(\vert E \vert)\) and depth \(O(D \cdot \log {d_\text{max}})\) where D is the graph diameter and \(d_\text{max}\) is the maximum vertex degree. This bound follows from the SpMSpV complexity described in Section 6.3.1. We are not aware of a worst-case analysis of the direction-optimized search but most switching heuristics are conservative and only switch to the bottom-up direction when it decreases work. Given that analysis in Section 6.3.2 shows that masked-SpMV depth is no larger than SpMSpV depth, we conclude that our implementation has worst-case work \(O({\it nnz}(\mathbf {A})) = O(\vert E \vert))\) and depth \(O(D \cdot \log d_\text{max})\).

7.2 Single-source Shortest-path

Given a source vertex \(s \in V\), a SSSP is a full exploration of weighted graph G that produces a distance array of all vertices v reachable from s, representing paths from s to each v such that the path distances are minimized.

• Hardwired GPU implementationCurrently the highest-performing SSSP algorithm implementation on the GPU is the work from Davidson et al. [26]. They provide two key optimizations in their SSSP implementation: (1) a load balanced graph traversal method, and (2) a priority queue implementation that reorganizes the workload.

• GraphBLAST implementationWe take a different approach from Davidson et al. to solve SSSP. We show that our approach both avoids the need for ad hoc data structures such as priority queues and wins in performance. The optimizations we use are: (1) generalized direction optimization, which is handled automatically within the mxv operation rather than inside the user’s graph algorithm code, and (2) sparsifying the set of active vertices after each iteration by comparing each active vertex to see whether or not it improved over the stored distance in the distance array. The second phase introduces two additional steps (compare Figures 10(b) and 10(e)). These facts make our SSSP implementation an adaptive variant of Bellman-Ford. While the worst-case work complexity is \(O(\vert E \vert \, \vert V \vert))\), it is much faster in practice due to (a) convergence being achieved significantly before \(\vert V \vert\) iterations, and (b) direction optimization. Each iteration has only \(O(\log {d_\text{max}})\) depth.

7.3 PageRank

The PageRank link analysis algorithm assigns a numerical weighting to each element of a hyperlinked set of documents, such as the World Wide Web, with the purpose of quantifying its relative importance within the set. The iterative method of computing PageRank gives each vertex an initial PageRank value and updates it based on the PageRank of its neighbors (this is the “pull” formulation of PageRank), until the PageRank value for each vertex converges. There are variants of the PageRank algorithm that stop computing PageRank for vertices that have converged already and also remove it from the set of active vertices. This is called adaptive PageRank [48] (also known as PageRankDelta). In this paper, we do not implement or compare against this variant of PageRank. We also acknowledge that different kinds of iterative solvers can be used for computing PageRank [39].

• Hardwired GPU implementationOne of the highest-performing implementations of PageRank is written by Khorasani, Vora, and Gupta [51]. In their system, they use solve the load imbalance and GPU underutilization problem with a GPU adoption of GraphChi’s Parallel Sliding Window scheme [53]. They call this preprocessing step “G-Shard” and combine it with a concatenated window method to group edges from the same source IDs. We realize that due to G-Shard’s preprocessing this comparison is not exactly fair to GraphBLAS, but include the comparison, because they are one of the leaders in PageRank performance and despite their preprocessing, our dynamic load-balancing is sufficient to make our implementation faster in the geomean (geometric mean).

• Gunrock implementationGunrock supports both pull- and push-based PageRank; its push-based implementation is in general faster than its pull-based implementation. For a fair algorithmic comparison, we measure against Gunrock’s pull-based implementation.

• GraphBLAST implementationIn GraphBLAST, we rely on the merge-based load-balancing scheme discussed in Section 6. The advantage of the merge-based scheme is that unlike Khorasani, Vora, and Gupta, we do not need any specialized storage format; the GPU is efficient enough to do the load-balancing on the fly. In terms of exploiting input sparsity, we demonstrate that our system is intelligent enough to determine that we are doing repeated matrix-vector multiplication and because the vector does not get any sparser, it is more efficient to use SpMV rather than SpMSpV. As described in Section 6.3.2, our SpMV implementation is based on ModernGPU [4], which uses the segmented-scan primitive with logarithmic depth and linear work. Consequently, an iteration of PageRank has \(O(\tilde{e})\) work and \(O(\lg {\tilde{e}})\) depth, where \(\tilde{e} = \sum _{i \vert \mathbf {m}(i) \ne 0} {{\it nnz}(\mathbf {A}(i,:))}\) is the number of nonzeros that need to be touched in that iteration. Note that \(\tilde{e}\) is at most \({\it nnz}(\mathbf {A})\) but is often smaller due to already converged vertices.

7.4 Connected Components

Weakly connected components (abbreviated as connected components) is the problem of: (1) identifying all subgraphs (or components) in an undirected graph such that every pair of vertices in the subgraph are connected by edges and no edges connect vertices in different subgraphs, and (2) labeling each vertex with its component ID.

• Hardwired GPU implementationSoman, Kishore and Narayanan [79] base their GPU implementation on two algorithms from the PRAM literature: hooking and pointer-jumping. Hooking takes an edge as the input and tries to set the component IDs of the two end vertices of that edge to the same value. In odd-numbered iterations, the lower vertex writes its value to the higher vertex, and vice versa in the even-numbered iteration. This strategy is used to increase the rate of convergence over a more naive approach such as a breadth-first-search. Pointer-jumping reduces a multi-level tree in the graph to a one-level tree (star). By repeating these two operators until no component ID changes for any node in the graph, the algorithm will compute the number of connected components for the graph and the connected component to which each node belongs.

• Gunrock implementationGunrock uses a filter operator on an edge frontier to implement hooking. The frontier starts with all edges and during each iteration, one end vertex of each edge in the frontier tries to assign its component ID to the other vertex, and the filter step removes the edge whose two end vertices have the same component ID. We repeat hooking until no vertex’s component ID changes and then proceed to pointer-jumping, where a filter operator on vertices assigns the component ID of each vertex to its parent’s component ID until it reaches the root. Then, a filter step removes the node whose component ID equals its own node ID. The pointer-jumping phase also ends when no vertex’s component ID changes.

• GraphBLAST implementationGraphBLAST’s implementation of CC is based on the FastSV algorithm [90]. FastSV is a linear-algebraic connected components algorithm [89] that is based on the classic PRAM algorithm of Shiloach and Vishkin [75] that uses hooking and pointer-jumping. We make two interesting observations. (1) We include a sparsification optimization that is discussed in the FastSV paper [90]. With three lines of code that zero out the redundant values, our push-pull design is able to handle this optimization automatically. (2) We avoid unnecessary GPU-to-CPU and CPU-to-GPU memory copies that would otherwise be required in the original FastSV GraphBLAS implementation using two new variants of assign and extract. Instead of using the Index * found with the standard variants, we use a GraphBLAS Vector that will implicitly have its values treated as the indices. Since Vector will always have the most recent data in GPU memory, this solves the problem of not being able to deduce whether the Index pointer is located in CPU or GPU memory. Unlike the other linear-algebraic CC implementation (LACC), FastSV does not guarantee \(O(\log {\vert V \vert })\) worst-case iteration bound because it avoids unconditional hooking for higher performance [89]. We chose FastSV to include in GraphBLAST because it is consistently faster than LACC in practice, despite having worse complexity bounds. Consequently, our CC implementation also takes worst-case \(O(\vert E \vert \, \vert V \vert)\) time and \(O(\vert V \vert)\) depth. However, it is much faster in practice, similar to other quadratic time and linear depth CC algorithms such as multistep-CC [78].

7.5 Triangle Counting

Triangle counting is the problem of counting the number of unique triplets \(u, v, w\) in an undirected graph such that \(\lbrace (u, v), (u,w), (v, w)\rbrace \in E\). Many important measures of a graph are triangle-based, such as clustering coefficient and transitivity ratio.

• Hardwired GPU implementationOne of the best-performing implementations of triangle counting is by Bisson and Fatica [12]. In their work, they demonstrate an effective use of a static workload mapping of thread, warp, block per matrix row together with using bitmaps.

• GraphBLAST implementationIn GraphBLAST, we follow Azad and Buluç [2] and Wolf et al. [83] in modeling the TC problem as a masked matrix-matrix multiplication problem. Given an adjacency matrix of an undirected graph \(\mathbf {A}\), and taking the lower triangular component \(\mathbf {L}\), the number of triangles is the reduction of the matrix \(\mathbf {B} = \mathbf {L}\mathbf {L}^T \,.\!* \mathbf {L}\) to a scalar. Rows of \(\mathbf {L}\) are sorted by increasing number of nonzeros, following the literature that demonstrates the benefits of sorting by degree prior to triangle counting [23]. In our implementation, we use a generalization of this algorithm where we assume we are solving the problem for three distinct matrices \(\mathbf {A}\), \(\mathbf {B}\), and \(\mathbf {M}\) by computing \(\mathbf {C}= \mathbf {A}\mathbf {B}\,.\!* \mathbf {M}\). We use the masked SpGEMM primitive whose implementation is detailed in Section 6.3.4. This is followed by a reduction of matrix \(\mathbf {C}\) to a scalar, returning the number of triangles in graph \(\mathbf {A}\). For each nonzero entry in the mask \(\mathbf {M}(i,j) \ne 0\), our masked matrix-matrix multiplication performs an intersection of the nonzeros in the i-th row of \(\mathbf {A}\) with the j-th column of \(\mathbf {B}\). Using a straightforward merge-based set intersection would have yielded an implementation with \(O(\vert E \vert ^{3/2})\) work and \(O(\lg ^{3/2} \vert V \vert)\) depth [77]. Instead of performing the set intersection using merging or hash tables, we use repeated binary searches from the elements of the shorter list to the larger list. Multiple publications concluded that the method of repeated binary searches was either competitive with or faster than the merge-based method on GPUs [36, 44]. Theoretically, it increases work marginally by a factor \(\log {d}\) on average where d is the average degree of a vertex. On the positive side, it has more parallelism.

8.1 Performance Summary

Table 12 and Figure 11 compare GraphBLAST’s performance against several other graph libraries and hardwired GPU implementations. In general, GraphBLAST’s performance on traversal-based algorithms (BFS and SSSP) is better on the seven scale-free graphs (soc-orkut, soc-lj, h09, i04, and rmats) than on the small-degree large-diameter graphs (rgg and road_usa). The main reason is our load-balancing strategy during traversal and particularly our emphasis on high performance for highly-skewed-distribution irregular graphs. Therefore, we incur a certain amount of overhead for our merge-based load-balancing and our requirement of a kernel launch in every iteration. For these types of graphs, asynchronous approaches, pioneered by Enterprise [55], that do not require exiting the kernel until the breakpoint has been met is a way to address the kernel launch problem. However, this does not work for non-BFS solutions, so asynchronous approaches in this area remain an open problem. In addition, graphs with uniformly low degree expose less parallelism and would tend to show smaller gains in comparison to CPU-based methods.

Fig. 11.

View Figure

Table 12.

View Table

Table 12. GraphBLAST’s Performance Comparison for Runtime and Edge Throughput with Other Graph Libraries (SuiteSparse, Ligra, Gunrock) and Hardwired GPU Implementations on a Tesla K40c GPU

8.2 Comparison with CPU Graph Frameworks

We compare GraphBLAST’s performance with three CPU graph libraries: the SuiteSparse GraphBLAS library, the first GraphBLAS implementation for multi-threaded CPU [27]; and Galois [62] and Ligra [76], both among the highest-performing multi-core shared-memory graph libraries. Against SuiteSparse, the speedup of GraphBLAST on average on all algorithms is geomean \(27.9\times\) (\(1268\times\) peak) and geomean \(43.51\times\) (\(1268\times\) peak) on scale-free graphs. Compared to Galois, GraphBLAST’s performance is generally faster. We are 2.6\(\times\) geomean (\(64.2\times\) peak) faster across all algorithms. We get the greatest speedup on BFS, because we implement direction optimization. We get the next greatest speedup on PR, where the amount of computation tends to be greater than for BFS or SSSP.

Compared to Ligra, GraphBLAST’s performance is generally comparable on most tested graph algorithms; note Ligra results are on a 2-CPU machine of the same timeframe as the K40c GPU we used to test. We are \(3.38\times\) (\(1.35\times\) peak) faster for BFS vs. Ligra for scale-free graphs, because we incorporate some BFS-specific optimizations such as masking, early-exit, and operand reuse, as discussed in Section 7. However, we are \(4.88\times\) slower on the road network graphs. For SSSP, a similar picture emerges. Compared to Ligra for scale-free graphs, we get \(1.35\times\) (\(1.72\times\) peak) speed-up, but are \(2.98\times\) slower on the road networks. We believe this is because our Bellman-Ford with sparsification means we can do less work on scale-free graphs, but our framework is not optimized for road networks. For PR, we are \(9.23\times\) (\(10.96\times\) peak) faster, because we use a highly-optimized merge-based load balancer that is suitable for this SpMV-based problem. For CC, we are \(1.30\times\) (\(1.88\times\) peak) faster. With regards to TC, we are \(2.80\times\) slower, because we have a simple algorithm for the masked matrix-matrix multiply.

8.3 Comparison with GPU Graph Frameworks and GPU Hardwired

Compared to hardwired GPU implementations, depending on the dataset, GraphBLAST’s performance is comparable or better on BFS, SSSP, and PR. For CC, GraphBLAST is \(3.1\times\) slower (geometric mean) on scale-free graphs and \(107.7\times\) slower on road network graphs. We think the reason is that we are strictly following Shiloach-Vishkin in doing only one level of pointer jumping per iteration, whereas the hardwired implementation is doing pointer jumping until each tree has been reduced to a star. This is an advantage on the GPU, because it allows their implementation to significantly reduce kernel launch overheads (see Section 2.1.4), which become significant especially for graphs with high diameter such as road networks. If kernel fusion is added to GraphBLAST, we would be able to take advantage of this optimization.

For TC, GraphBLAST is \(3.3\times\) slower (geometric mean) than the hardwired GPU implementation due to fusing of the matrix-multiply and the reduce, which lets the hardwired implementation avoid the step of writing out the output to the matrix-multiply. The alternative is having a specialized kernel that does a fused matrix-multiply and reduce. This tradeoff is not typical of our other algorithms. While still achieving high performance, GraphBLAST’s application code is smaller in size and clearer in logic compared to other GPU graph libraries.

Compared to CuSha and MapGraph, GraphBLAST’s performance is quite a bit faster. We get geomean speedups of \(8.40\times\) and \(3.97\times\) respectively (\(420\times\) and \(64.2\times\) peak). The speedup comes from direction optimization. CuSha only does the equivalent of pull-traversal, so their performance is most comparable to ours in PR. MapGraph is push-only.

Compared to Gunrock, the fastest GPU graph framework, GraphBLAST’s performance is comparable on BFS, CC, and TC with Gunrock being 11.8%, 14.8%, and 11.1% faster in the geomean, respectively. On SSSP, GraphBLAST is faster by \(1.1\times\) (\(1.53\times\) peak). This can be attributed to GraphBLAST using generalized direction optimization and Gunrock only doing push-based advance. On PR, GraphBLAST is significantly faster and gets speedups of \(2.39\times\) (\(5.24\times\) peak). For PR, the speed-up again can be attributed to GraphBLAST automatically using generalized direction optimization to select the right direction, which is SpMV in this case. Gunrock does push-based advance.

8.4 Comparison with Gunrock on Latest GPU Architecture

In Table 13, we compare against Gunrock on BFS, SSSP, CC, and PR using the latest generation GPU, Titan V. As the result shows, we have a 3.13\(\times\) slowdown compared to Gunrock on BFS, which indicates that we do worse on Titan V. On SSSP, we are \(1.08\times\) (\(2.39\times\) peak) faster when not including the road network datasets, and \(0.48\times\) slower when including them. On CC, we have a 4.00\(\times\) slowdown compared to Gunrock. On PR, we are \(2.90\times\) (\(7.67\times\) peak) faster in the geomean.

Table 13.

View Table

Table 13. GraphBLAST’s Performance Comparison for Runtime and Edge Throughput with Gunrock [82] for Four Graph Algorithms on a Titan V GPU

Taking a closer look at this in Figure 12, we can see that both push and pull components of Gunrock’s BFS benefit due to moving from K40c to Titan V, but “other” does not. However for GraphBLAST, only the “other” and pull benefit from moving from K40c to Titan V. We hypothesize the reason for this is the GraphBLAST push is implemented using a radixsort to perform a multiway merge, and radixsort does not see a noticeable improvement in performance from K40c to Titan V on the problem sizes typical of BFS. On the other hand, Gunrock uses uses a series of inexpensive heuristics [82] to reduce but not eliminate redundant entries in the output frontier. These heuristics include a global bitmask, a block-level history hashtable, and a warp-level hashtable. The size of each hashtable is adjustable to achieve the optimal tradeoff between performance and redundancy reduction rate. However, this approach may not be suitable for GraphBLAS, because such an optimization may be too BFS-focused and would generalize poorly.

Fig. 12.

View Figure

For SSSP, GraphBLAST has the advantage of the direction optimization automatically choosing the optimal direction, which for SSSP happens to be pull. Gunrock’s SSSP is written to use push, so it is unable to take advantage of this feature. For PR, the situation is similar to SSSP in that Gunrock uses push when pull is better. In CC, Gunrock’s implementation has the advantage of following the approach of the hardwired implementation by Soman, Kishore and Narayanan [79]. This algorithm performs pointer-jumping until the tree has become a star, whereas GraphBLAST follows Shiloach-Vishkin strictly and only does one level of pointer-jumping. We believe Gunrock’s primary performance advantage, then, is faster convergence. More research is required to understand whether this optimization can be used to improve the linear algebraic implementation of FastSV in the presence and absence of kernel fusion (see Section 2.1.4).

8.5 Comparison with GraphBLAS-like Framework on GPU

In Table 14, we compare against GBTL [88], the first GraphBLAS-like implementation for the GPU on BFS. Our implementation is \(31.8\times\) (\(58.5\times\) peak) faster in the geomean. We attribute this speed-up to several factors: (1) they use the Thrust library [10] to manage the CPU-to-GPU memory traffic that works for all GPU applications, while we use a domain-specific memory allocator that only copies from CPU to GPU when necessary; (2) they specialize the CUSP library’s mxm operation [24, 25] for a matrix with a single column to mimic the mxv required by the BFS, while we have a specialized mxv operation that is more efficient; and (3) we utilize the design principles of exploiting input and output sparsity, as well as proper load balancing, none of which are in GBTL.

In addition to getting comparable or faster performance, GraphBLAST has the advantage of being concise, as shown in Table 1. Developing new graph algorithms in GraphBLAST requires modifying a single file and writing straightforward C++ code. Currently, we are working on a Python frontend interface too, to allow users to build new graph algorithms without having to recompile. Additional language bindings are being planned as well (see Figure 13). Similar to working with machine learning frameworks, writing GraphBLAST code does not require any parallel programming knowledge of OpenMP, OpenCL or CUDA, or even performance optimization experience.

Fig. 13.

View Figure