A GPU-accelerated reachability query method and system

Abstract

This invention discloses a GPU-accelerated reachability query method, comprising: receiving graph data from a user; performing compression preprocessing on the graph data to obtain compressed graph data; obtaining query points from the user, including source vertices and target vertices; configuring the GPU environment; copying the compressed graph data and query points into the GPU memory within the configured GPU environment; calculating the average out-degree of the graph data and determining whether the average out-degree is greater than a pre-set threshold; if so, obtaining the graph data copied to the GPU memory; using a two-stage BFS algorithm to query whether the obtained query points are reachable in the graph data as the query result; and then copying the query result to CPU memory and outputting it. This invention solves the technical problem that existing index-based reachability query algorithms cannot handle large-scale graph data because building an index for reachability queries requires a large amount of storage resources.

Description

Technical Field

[0001] This invention belongs to the field of graph computing technology, and more specifically, relates to a GPU-accelerated reachability query method and system. Background Technology

[0002] The rapid development of the internet and the advent of the big data era have led to an increasing demand for the rapid analysis of massive amounts of data. Graphs are a general and effective data structure for representing various types of objects. Vertices in a graph describe entities, while edges describe the relationships between entities. For example, in social networking platforms, users are considered nodes, and interactions between users are represented by edges, effectively depicting the connection patterns of the social network. Furthermore, many non-graph structures and applications can be transformed into graph models for processing and analysis, including protein-protein interaction networks, gene regulatory networks, knowledge graphs, and hardware architecture design. Due to its wide range of applications, graph data processing research has become a key area of ​​focus for both academia and industry.

[0003] Graph computing is the field that studies how to accelerate the computation, storage, and management of large-scale graph data. All analytical computations based on graph data fall under graph computing, thus its applications are extremely broad. As a classic problem in graph computing, reachability query (RQ) is a research hotspot in computer science, geographic information science, and transportation, and can be used to solve optimization problems such as resource allocation, network connectivity analysis, and path planning. Furthermore, reachability query is a fundamental technique for problems such as community queries in graph theory.

[0004] There is currently a large amount of research on reachability query technology, including index-based reachability query and distributed reachability query. Among them, index-based reachability query achieves fast querying of reachability between nodes by pre-constructing an index structure of reachability information for graph data; distributed reachability query distributes large-scale graph data across multiple machines and processes queries in parallel through a distributed computing framework.

[0005] However, both of the above-mentioned existing reachability query methods have some significant drawbacks: First, index-based reachability query algorithms incur substantial storage overhead when maintaining a large number of index data structures, making them unsuitable for handling large-scale graph data; Second, distributed reachability query algorithms can lead to load imbalance when sharding large-scale graph data, and frequent data exchanges between nodes can result in significant communication overhead. Summary of the Invention

[0006] In view of the above-mentioned defects or improvement needs of the existing technology, the present invention provides a GPU-accelerated reachability query method and system. Its purpose is to solve the technical problems of existing index-based reachability query algorithms, which cannot handle large-scale graph data due to the large amount of storage resources required to build indexes for reachability queries, and the technical problems of existing distributed reachability query algorithms, which cause unbalanced load and high communication overhead when performing data sharding on large-scale graph data.

[0007] To achieve the above objectives, according to one aspect of the present invention, a GPU-accelerated reachability query method is provided, comprising the following steps:

[0008] (1) Receive graph data from the user and perform compression preprocessing on the graph data to obtain compressed graph data.

[0009] (2) Obtain the query points from the user, including the source vertex and the target vertex, configure the GPU environment, and copy the graph data and query points after compression in step (1) into the GPU memory in the configured GPU environment;

[0010] (3) Calculate the average out-degree of the graph data copied to the GPU memory in step (2), and determine whether the average out-degree is greater than the preset threshold (the value range is 3 to 9, preferably 8). If it is, proceed to step (4); otherwise, proceed to step (5).

[0011] (4) Obtain the graph data copied to the GPU memory in step (2), and use the two-stage BFS algorithm to query whether the query point obtained in step (2) is reachable in the graph data as the query result, and then proceed to step (6).

[0012] (5) Obtain the graph data copied to the GPU memory in step (2), and use the bidirectional BFS algorithm to query whether the query point obtained in step (2) is reachable in the graph data as the query result, and then proceed to step (6).

[0013] (6) Copy the query results to CPU memory and output them.

[0014] Preferably, step (1) specifically includes the following sub-steps:

[0015] (1-1) Receive graph data from the user, set counter i = 0, obtain the total number of edges and vertices in the graph data, and initialize array C and array R to empty;

[0016] (1-2) Determine if i is equal to the total number of vertices in the graph data. If yes, use the obtained array C as the final adjacency list and array R as the final offset list. Array C and array R are the compressed graph data. The process ends. Otherwise, proceed to step (1-3).

[0017] (1-3) Determine whether there is an adjacent vertex of the i-th vertex in the graph data. If so, add the adjacent vertices to array C in sequence, and record the position of the first adjacent vertex of the i-th vertex in the graph data in array C in array R. Then proceed to step (1-5). Otherwise, proceed to step (1-4).

[0018] (1-4) Set the i-th element R[i] in array R to the (i-1)-th element R[i-1], and then proceed to step (1-5);

[0019] (1-5) Set i = i + 1 and return to step (1-2).

[0020] Preferably, the average out-degree σ in step (3) is equal to:

[0021]

[0022] Where n represents the total number of vertices in the graph data, and out_deg(i) represents the out-degree of the i-th vertex in the graph data.

[0023] Preferably, this step specifically includes the following sub-steps:

[0024] (4-1) Create a vertex queue vque, an edge queue eque, and a distance array dis. Initialize the vertex queue vque to the source vertex, initialize the edge queue eque to empty, and initialize the distances corresponding to the source vertex in the distance array dis to 0, and initialize the distances corresponding to the other vertices in the distance array dis to infinity.

[0025] (4-2) Determine if the vertex queue vque is empty. If it is, take the source vertex and the target vertex as unreachable as the query result and the process ends. Otherwise, determine if the vertex queue vque contains the target vertex. If it is, take the source vertex and the target vertex as reachable as the query result and the process ends. Otherwise, proceed to step (4-3).

[0026] (4-3) Get the global ID of the current thread;

[0027] (4-4) Set the counter global_id1 to the global ID of the current thread obtained in step (4-3);

[0028] (4-5) Determine if global_id1 is less than the length of the vertex queue vque. If it is, get the global_id1-th vertex vque[global-id1] in the vertex queue vque and then proceed to step (4-6). Otherwise, it means that all vertices in the vertex queue vque have been processed and then proceed to step (4-9).

[0029] (4-6) Get the size of the adjacency list corresponding to the global_id1 vertex vque[global-id1] in the vertex queue vque. Based on the size of the adjacency list, get the prefix sum of the thread bundle corresponding to the thread of the vertex in the thread grid. Use this as the offset of the adjacent vertices of the global_id1 vertex vque[global-id1] in the corresponding adjacency list in the edge queue eque. Store all adjacent vertices in the edge queue eque.

[0030] (4-7) Set global_id1 = global_id1 + gridDim.x * blockDim.x, and return to step (4-5); where gridDim.x represents the number of thread blocks in the thread grid where the thread of dimension x is located, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and gridDim.x * blockDim.x represents the total number of threads in the thread grid where the thread of dimension x is located;

[0031] (4-8) Copy the edge queue eque obtained in step (4-6) to the shared memory, and use a hash table in the shared memory to perform collision detection processing on the edge queue eque (i.e. remove duplicate vertices) to obtain the edge queue eque after collision detection processing.

[0032] (4-9) Set the counter global_id2 to the global ID of the current thread obtained in step (4-3);

[0033] (4-10) Determine if global_id2 is less than the length of the edge queue eque. If it is, get the global_id2-th vertex eque[global_id2] in the edge queue eque and then proceed to step (4-11). Otherwise, it means that all vertices in the edge queue eque have been processed and then return to step (4-2).

[0034] (4-11) Determine whether the distance to the second global_id2 vertex eque[global_id2] in the distance array dis is infinite. If it is, update the distance to the second global_id2 vertex eque[global_id2] in the distance array dis and add the second global_id2 vertex eque[global_id2] to the vertex queue vque. Otherwise, proceed to step (4-12).

[0035] (4-12) Set global_id2 = global_id2 + gridDim.x * blockDim.x and return to step (4-10).

[0036] Preferably, step (4-6) specifically includes the following sub-steps:

[0037] (4-6-1) Get the size of the adjacency list corresponding to the global_id1 vertex vque[global-id1] in the vertex queue vque, set the counter i=0, initialize the array prefix to empty, which is used to record the prefix sum result of the thread corresponding to the global_id1 vertex in the thread grid, initialize the array nbr (its size is 32) to empty, which is used to record the size of the adjacency list corresponding to the global_id1 vertex, and initialize the 0th element prefix[0] in the array prefix to 0;

[0038] (4-6-2) Determine if i is equal to the size of array prefix - 1. If yes, use array prefix as the final prefix sum result table and the process ends. Otherwise, proceed to step (4-6-3).

[0039] (4-6-3) Set the i-th element in the array prefix to prefix[i-1]+nbr[i-1], where nbr[i-1] represents the size of the adjacency list corresponding to the global_id1 vertex, and then proceed to step (4-6-4);

[0040] (4-6-4) Set i = i + 1 and return to step (4-6-2).

[0041] Preferably, in step (4-3), the global ID of the thread is equal to:

[0042] ID=blockIdx.x*blockDim.x+threadIdx.x

[0043] Where blockIdx.x represents the number of the thread block corresponding to the thread of dimension x in the entire thread grid, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and threadIdx.x represents the number of the thread bundle in the thread grid in which the thread of dimension x is located.

[0044] Preferably, step (5) specifically includes the following sub-steps:

[0045] (5-1) Create a forward queue fque, a backward queue bque, a forward distance array fdis, and a backward distance array bdis. Initialize the forward queue fque as the source vertex, the backward queue bque as the target vertex, initialize the distance corresponding to the source vertex in the forward distance array fdis to 0, and initialize the distance corresponding to the other vertices to infinity. Initialize the distance corresponding to the target vertex in the backward distance array bdis to 0, and initialize the distance corresponding to the other vertices to infinity.

[0046] (5-2) Determine whether there are distances in the forward distance array fdis and the backward distance array bdis that are not infinite at the same position. If they are, the source vertex and the target vertex are reachable as the query result and the process ends. Otherwise, determine whether the forward queue fque and the backward queue bque are both empty. If they are, the source vertex and the target vertex are not reachable as the query result and the process ends. Otherwise, proceed to step (5-3).

[0047] (5-3) Get the global ID of the current thread;

[0048] (5-4) Set the counter global_id3 to the global ID of the current thread obtained in step (5-3);

[0049] (5-5) Determine if global_id3 is less than the length of the forward queue fque. If it is, get the 3rd global_id3 vertex fque[global_id3] in the forward queue fque and then proceed to step (5-6). Otherwise, it means that all vertices in the forward queue fque have been processed and then proceed to step (5-8).

[0050] (5-6) Determine whether the distance between the adjacent vertices in the adjacency list corresponding to the 3rd global_id vertex fque[global_id3] in the forward queue fque is infinite. If so, update the distance between the adjacent vertices in the adjacency list corresponding to the 3rd global_id vertex fque[global_id3] in the forward queue fque and the distance in the forward distance array fdis, and add the adjacent vertices in the adjacency list corresponding to the 3rd global_id vertex fque[global_id3] in the forward queue fque. Otherwise, it means that the 3rd global_id vertex fque[global_id3] in the forward queue fque has been processed. Remove the vertex from the forward queue fque and then proceed to step (5-7).

[0051] (5-7) Set global_id3 = global_id3 + gridDim.x * blockDim.x and return to step (5-5); where gridDim.x represents the number of thread blocks in the thread grid where the thread of dimension x is located, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and gridDim.x * blockDim.x represents the total number of threads in the thread grid where the thread of dimension x is located (if the number of elements in the vertex queue is much larger than the total number of threads (gridDim.x * blockDim.x), this grid-striding loop mode can enable each thread to process multiple data elements and make reasonable use of computing resources).

[0052] (5-8) Set the counter global_id4 to the global ID of the current thread obtained in step (5-3);

[0053] (5-9) Determine if global_id4 is less than the length of the back queue bque. If it is, get the 4th vertex bque[global_id4] in the back queue bque and then proceed to step (5-10). Otherwise, it means that all vertices in the back queue bque have been processed and then return to step (5-2).

[0054] (5-10) Determine whether the distance between the adjacent vertices in the adjacency list corresponding to the 4th global_id vertex bque[global_id4] in the backward queue bque is infinite. If so, update the distance between the adjacent vertices in the adjacency list corresponding to the 4th global_id vertex bque[global_id4] in the backward queue bque and the distance between the adjacent vertices in the backward distance array bdis, and add the adjacent vertices in the adjacency list corresponding to the 4th global_id vertex bque[global_id4] in the backward queue bque. Otherwise, it means that the 4th global_id vertex bque[global_id4] in the backward queue bque has been processed. Remove the vertex from the backward queue bque, and then proceed to step (5-11).

[0055] (5-11) Set global_id4 = global_id4 + gridDim.x * blockDim.x and return to step (5-9).

[0056] According to another aspect of the present invention, a GPU-accelerated reachability query system is provided, comprising:

[0057] The first module is used to receive graph data from the user, compress and preprocess the graph data to obtain compressed graph data.

[0058] The second module is used to obtain query points from users, including source vertices and target vertices, configure the GPU environment, and copy the graph data compressed by the first module and the query points into the GPU memory in the configured GPU environment.

[0059] The third module is used to calculate the average out-degree of the graph data copied to the GPU memory by the second module, and to determine whether the average out-degree is greater than a preset threshold (its value range is 3 to 9, preferably 8). If it is, it proceeds to the fourth module; otherwise, it proceeds to the fifth module.

[0060] The fourth module is used to obtain the graph data copied to the GPU memory by the second module, and use a two-stage BFS algorithm to query whether the query point obtained by the second module is reachable in the graph data as the query result, and then the sixth module;

[0061] The fifth module is used to obtain the graph data copied to the GPU memory by the second module, and use the bidirectional BFS algorithm to query whether the query point obtained by the second module is reachable in the graph data as the query result, and then transfer to the sixth module.

[0062] The sixth module is used to copy the query results to CPU memory and output them.

[0063] In summary, compared with the prior art, the above-described technical solutions conceived by this invention can achieve the following beneficial effects:

[0064] (1) Since the present invention adopts steps (1-1) to (1-5), it mainly achieves compressed storage of large-scale graph data by storing the adjacency list and offset information of the graph data. Therefore, it can solve the problem that the existing index-based reachability query algorithm cannot handle large-scale graph data due to the large storage space occupied by the index.

[0065] (2) Since the present invention adopts steps (4-6), it mainly provides accurate workload statistics and dynamic work allocation mechanism by calculating prefix sum, making parallel processing more effective. Therefore, it can solve the problem of unbalanced load caused by data sharding in the existing distributed reachability query algorithm.

[0066] (3) Since the present invention adopts steps (4-9) to (4-13), it mainly reduces data access latency by performing information interaction in shared memory, making communication and synchronization faster. Therefore, it can solve the problem that existing distributed reachability query algorithms will have a large communication overhead due to frequent data exchange between nodes.

[0067] (4) By using step (3), the present invention employs different algorithms for graphs with different average out-degrees. This can not only reasonably allocate tasks in parallel computing and improve resource utilization, but also select different data structures to reduce unnecessary calculations and speed up querying. Attached Figure Description

[0068] Figure 1 This is a schematic diagram of a typical device memory model in CUDA;

[0069] Figure 2 This is a schematic diagram illustrating the principle of CUDA parallel computing;

[0070] Figure 3 This is a schematic diagram of compressed storage in this invention;

[0071] Figure 4 This is a schematic diagram of prefix sum calculation in this invention;

[0072] Figure 5 This is a flowchart of the GPU-accelerated reachability query method of the present invention. Detailed Implementation

[0073] To make the objectives, technical solutions, and advantages of this invention clearer, the invention will be further described in detail below with reference to the accompanying drawings and embodiments. It should be understood that the specific embodiments described herein are merely illustrative and not intended to limit the invention. Furthermore, the technical features involved in the various embodiments of this invention described below can be combined with each other as long as they do not conflict with each other.

[0074] The technical terms used in this invention will be explained and described below:

[0075] Reachability Query: Given a graph G = (V, E), where V is the set of vertices and E is the set of edges, for any two vertices s and t in G, the reachability query is denoted by q(s, t). If there exists at least one path from s to t in G, then s is said to be reachable from t, denoted as q(s, t) = true. If there is no path from s to t, then s is said to be unreachable from t, denoted as q(s, t) = false.

[0076] Parallel computing refers to solving a problem by using multiple processing units (such as CPU cores, GPU cores, and compute nodes) simultaneously, thereby accelerating the computation process. The core idea is to decompose the computational task into multiple subtasks that can be executed in parallel, and then use multiple processing units to execute these tasks simultaneously, thus improving efficiency and speed.

[0077] Thread: The basic unit for executing computational tasks. A GPU can execute a large number of threads simultaneously, which can execute the same instructions in parallel, thereby accelerating the computation process.

[0078] A warp is the smallest unit of hardware scheduling in an NVIDIA GPU, typically containing 32 threads. All threads in a warp share the same program counter, meaning they execute the same instructions at the same time.

[0079] Thread Block: A fundamental concept in the CUDA programming model of NVIDIA GPU architecture. It is a unit of organization for threads, used to perform parallel computing tasks on the GPU. Each thread block contains multiple thread bundles, and a thread block is a basic scheduling unit that can be executed in parallel on the GPU.

[0080] Thread Grid: A thread grid is a combination of multiple thread blocks. Thread grids provide a way to organize and manage large numbers of thread blocks, enabling programs to execute efficiently in parallel on the GPU.

[0081] Shared memory is a storage area that is exclusively used by each thread block, and all threads within the same thread block can access and manipulate this memory. A major advantage of shared memory is that its access latency is lower than that of global memory, making it suitable for collaboration and data sharing between threads within a thread block.

[0082] Figure 1 This is a typical framework for the device memory model in Compute Unified Device Architecture (CUDA). It includes: first, copying data from the host (CPU) to the global memory of the device (GPU); then, performing thread-level parallel computation on the device, distributing a large number of computational tasks to different threads in a Single Instruction Multiple Data (SIMD) manner; the grid is the highest-level thread organization unit, consisting of multiple thread blocks; threads within the same thread block can choose to store intermediate results in shared memory with faster access speed, which can significantly improve computation speed; and finally, copying the computation results from the device (GPU) back to the host (CPU).

[0083] Figure 2 Based on the principles of CUDA parallel computing, in this invention, a thread block contains 256 threads. Thread global IDs are indexed using the thread block number (blockIdx.x) and the thread number (threadIdx.x). All threads process their respective tasks based on their IDs. The GPU internally includes a large amount of global memory and a small amount of shared memory. Threads within a thread block communicate quickly through shared memory and then write the results to global memory for use by other thread blocks.

[0084] The basic idea of ​​this invention is to achieve kernel-level task parallelism through an optimized GPU-accelerated reachability query method. This invention first compresses the graph data to obtain graph data that occupies less storage space, and then copies the compressed graph data to GPU memory for querying. Graph data with a large average out-degree is queried using a two-phase Breadth-First Search (BFS) algorithm, while graph data with a small average out-degree is queried using a bidirectional BFS algorithm. By using different algorithms based on the graph data structure and resource status and starting an appropriate number of threads, the query processing can be executed in parallel.

[0085] like Figure 5 As shown, this invention provides a GPU-accelerated reachability query method, comprising the following steps:

[0086] (1) Receive graph data from the user and perform compression preprocessing on the graph data to obtain compressed graph data.

[0087] The advantage of this step is that using compressed graph data allows large-scale graph data to occupy less storage space, speeds up data transfer between the host and the device, enables the GPU to process larger-scale graph data, and the compressed data structure is more suitable for parallel processing.

[0088] like Figure 3 As shown, this step specifically includes the following sub-steps:

[0089] (1-1) Receive graph data from the user, set counter i = 0, and obtain the total number of edges and vertices in the graph data, and initialize array C (whose size is equal to the total number of edges) and array R (whose size is equal to the total number of vertices) to be empty;

[0090] Specifically, array C is used to record the adjacent vertices of all vertices in the graph data, and array R is used to record the offset of the starting position of the adjacent vertices of all vertices in the graph data within array C.

[0091] More specifically, the array is initialized in this step by obtaining it from the computer using input commands, for example, by using the C++ statement "vector<int>". <int>C".

[0092] (1-2) Determine if i is equal to the total number of vertices in the graph data. If yes, use the obtained array C as the final adjacency list and array R as the final offset list. Array C and array R are the compressed graph data. The process ends. Otherwise, proceed to step (1-3).

[0093] (1-3) Determine whether there is an adjacent vertex of the i-th vertex in the graph data. If so, add the adjacent vertices to array C in sequence, and record the position of the first adjacent vertex of the i-th vertex in the graph data in array C in array R. Then proceed to step (1-5). Otherwise, proceed to step (1-4).

[0094] (1-4) Set the i-th element R[i] in array R to the (i-1)-th element R[i-1], and then proceed to step (1-5);

[0095] (1-5) Set i = i + 1 and return to step (1-2).

[0096] (2) Obtain the query points from the user, including the source vertex and the target vertex, configure the GPU environment, and copy the graph data and query points after compression in step (1) into the GPU memory in the configured GPU environment;

[0097] Specifically, configuring the GPU environment in this step involves installing the NVIDIA driver and CUDA toolkit for the GPU, and installing the nvcc compiler.

[0098] (3) Calculate the average out-degree of the graph data copied to the GPU memory in step (2), and determine whether the average out-degree is greater than the preset threshold (the value range is 3 to 9, preferably 8). If it is, proceed to step (4); otherwise, proceed to step (5).

[0099] Specifically, in graph data, degree is divided into in-degree and out-degree, which represent the number of edges pointing to nodes in the graph data and the number of edges originating from that node, respectively.

[0100] In this step, the average out-degree σ is equal to:

[0101]

[0102] Where n represents the total number of vertices in the graph data, and out_deg(i) represents the out-degree of the i-th vertex in the graph data.

[0103] If the average out-degree σ exceeds the preset threshold, it indicates that the graph data structure is relatively dense, resulting in many intermediate results during the expansion process, and a high probability that vertex deduplication will be required later. In this case, a two-stage Breadth First Search (BFS) algorithm is recommended. Conversely, if the average out-degree σ does not exceed the threshold, it indicates that the graph data structure is relatively sparse, resulting in fewer intermediate results during the expansion process, and a lower probability that vertex deduplication will be required later. In this case, a bidirectional BFS algorithm is recommended.

[0104] (4) Obtain the graph data copied to the GPU memory in step (2), and use the two-stage BFS algorithm to query whether the query point obtained in step (2) is reachable in the graph data as the query result, and then proceed to step (6).

[0105] The advantage of this step is that it uses the collision detection feature of hash tables within shared memory to deduplicate intermediate results, and uses prefix sums and atomic operations to ensure accuracy and concurrency safety when multiple threads write data simultaneously, thus ensuring full utilization of computing resources.

[0106] like Figure 5 As shown, this step specifically includes the following sub-steps:

[0107] (4-1) Create a vertex queue vque, an edge queue eque, and a distance array dis. Initialize the vertex queue vque to the source vertex, initialize the edge queue eque to empty, and initialize the distances corresponding to the source vertex in the distance array dis to 0, and initialize the distances corresponding to the other vertices in the distance array dis to infinity.

[0108] (4-2) Determine if the vertex queue vque is empty. If it is, take the source vertex and the target vertex as unreachable as the query result and the process ends. Otherwise, determine if the vertex queue vque contains the target vertex. If it is, take the source vertex and the target vertex as reachable as the query result and the process ends. Otherwise, proceed to step (4-3).

[0109] (4-3) Get the global ID of the current thread;

[0110] Specifically, each thread uses the global ID as an initial index to process the vertices in the queue;

[0111] More specifically, the global ID of a thread is equal to:

[0112] ID=blockIdx.x*blockDim.x+threadIdx.x

[0113] Where blockIdx.x represents the number of the thread block corresponding to the thread of dimension x in the entire thread grid, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and threadIdx.x represents the number of the thread of dimension x in the thread bundle in its own thread grid. By multiplying the first two and adding the third, the global ID of the thread in the entire thread grid can be calculated.

[0114] (4-4) Set the counter global_id1 to the global ID of the current thread obtained in step (4-3);

[0115] (4-5) Determine if global_id1 is less than the length of the vertex queue vque. If it is, get the global_id1-th vertex vque[global-id1] in the vertex queue vque and then proceed to step (4-6). Otherwise, it means that all vertices in the vertex queue vque have been processed and then proceed to step (4-9).

[0116] (4-6) Get the size of the adjacency list corresponding to the global_id1 vertex vque[global-id1] in the vertex queue vque. Based on the size of the adjacency list, get the prefix sum of the thread bundle corresponding to the thread of the vertex in the thread grid. Use this as the offset of the adjacent vertices of the global_id1 vertex vque[global-id1] in the corresponding adjacency list in the edge queue eque. Store all adjacent vertices in the edge queue eque.

[0117] like Figure 4 As shown, this step specifically includes the following sub-steps:

[0118] (4-6-1) Get the size of the adjacency list corresponding to the global_id1 vertex vque[global-id1] in the vertex queue vque, set the counter i=0, initialize the array prefix (its size is 32) to be empty, which is used to record the prefix sum result of the thread corresponding to the global_id1 vertex in the thread grid it is in, initialize the array nbr (its size is 32) to be empty, which is used to record the size of the adjacency list corresponding to the global_id1 vertex, and initialize the 0th element prefix[0] in the array prefix to 0;

[0119] (4-6-2) Determine if i is equal to the size of array prefix - 1. If yes, use array prefix as the final prefix sum result table and the process ends. Otherwise, proceed to step (4-6-3).

[0120] (4-6-3) Set the i-th element in the array prefix to prefix[i-1]+nbr[i-1], where nbr[i-1] represents the size of the adjacency list corresponding to the global_id1 vertex, and then proceed to step (4-6-4);

[0121] (4-6-4) Set i = i + 1 and return to step (4-6-2).

[0122] The advantage of steps (4-6-1) to (4-6-4) above is that the prefix sum result is used to determine the offset of each thread in the process of writing to the side queue, so that multiple threads can write intermediate results to the side queue in parallel, reasonably allocate the side queue space, reduce thread contention, and achieve better load balancing.

[0123] (4-7) Set global_id1 = global_id1 + gridDim.x * blockDim.x, and return to step (4-5); where gridDim.x represents the number of thread blocks in the thread grid where the thread of dimension x is located, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and gridDim.x * blockDim.x represents the total number of threads in the thread grid where the thread of dimension x is located;

[0124] If the number of elements in the vertex queue (vque) is much larger than the total number of threads (gridDim.x * blockDim.x), this grid-striding loop mode allows each thread to process multiple data elements, making efficient use of computing resources.

[0125] (4-8) Copy the edge queue eque obtained in step (4-6) to the shared memory, and use a hash table in the shared memory to perform collision detection processing on the edge queue eque (i.e. remove duplicate vertices) to obtain the edge queue eque after collision detection processing.

[0126] (4-9) Set the counter global_id2 to the global ID of the current thread obtained in step (4-3);

[0127] (4-10) Determine if global_id2 is less than the length of the edge queue eque. If it is, get the global_id2-th vertex eque[global_id2] in the edge queue eque and then proceed to step (4-11). Otherwise, it means that all vertices in the edge queue eque have been processed and then return to step (4-2).

[0128] (4-11) Determine whether the distance to the second global_id2 vertex eque[global_id2] in the distance array dis is infinite. If it is, update the distance to the second global_id2 vertex eque[global_id2] in the distance array dis and add the second global_id2 vertex eque[global_id2] to the vertex queue vque. Otherwise, proceed to step (4-12).

[0129] (4-12) Set global_id2 = global_id2 + gridDim.x * blockDim.x, and return to step (4-10);

[0130] (5) Obtain the graph data copied to the GPU memory in step (2), and use the bidirectional BFS algorithm to query whether the query point obtained in step (2) is reachable in the graph data as the query result, and then proceed to step (6).

[0131] The advantage of this step is that it uses graph data with a low average out-degree, which produces fewer intermediate results during the expansion process. Therefore, it can take advantage of the high concurrency of the GPU to process nodes in both search directions at the same time, increase the number of intermediate results, make full use of GPU computing and storage resources, and accelerate the query process.

[0132] like Figure 5 As shown, this step specifically includes the following sub-steps:

[0133] (5-1) Create a forward queue fque, a backward queue bque, a forward distance array fdis, and a backward distance array bdis. Initialize the forward queue fque as the source vertex, the backward queue bque as the target vertex, initialize the distance corresponding to the source vertex in the forward distance array fdis to 0, and initialize the distance corresponding to the other vertices to infinity. Initialize the distance corresponding to the target vertex in the backward distance array bdis to 0, and initialize the distance corresponding to the other vertices to infinity.

[0134] (5-2) Determine whether there are distances in the forward distance array fdis and the backward distance array bdis that are not infinite at the same position. If they are, the source vertex and the target vertex are reachable as the query result and the process ends. Otherwise, determine whether the forward queue fque and the backward queue bque are both empty. If they are, the source vertex and the target vertex are not reachable as the query result and the process ends. Otherwise, proceed to step (5-3).

[0135] (5-3) Get the global ID of the current thread;

[0136] Specifically, each thread uses the global ID as the initial index to process the vertices in queue e;

[0137] More specifically, the thread global ID is equal to:

[0138] ID=blockIdx.x*blockDim.x+threadIdx.x

[0139] Where blockIdx.x represents the number of the thread block corresponding to the thread of dimension x in the entire thread grid, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and threadIdx.x represents the number of the thread of dimension x in the thread bundle in its own thread grid. By multiplying the first two and adding the third, the global ID of the thread in the entire thread grid can be calculated.

[0140] (5-4) Set the counter global_id3 to the global ID of the current thread obtained in step (5-3);

[0141] (5-5) Determine if global_id3 is less than the length of the forward queue fque. If it is, get the 3rd global_id3 vertex fque[global_id3] in the forward queue fque and then proceed to step (5-6). Otherwise, it means that all vertices in the forward queue fque have been processed and then proceed to step (5-8).

[0142] (5-6) Determine whether the distance between the adjacent vertices in the adjacency list corresponding to the 3rd global_id vertex fque[global_id3] in the forward queue fque is infinite. If so, update the distance between the adjacent vertices in the adjacency list corresponding to the 3rd global_id vertex fque[global_id3] in the forward queue fque and the distance in the forward distance array fdis, and add the adjacent vertices in the adjacency list corresponding to the 3rd global_id vertex fque[global_id3] in the forward queue fque. Otherwise, it means that the 3rd global_id vertex fque[global_id3] in the forward queue fque has been processed. Remove the vertex from the forward queue fque and then proceed to step (5-7).

[0143] (5-7) Set global_id3 = global_id3 + gridDim.x * blockDim.x and return to step (5-5); where gridDim.x represents the number of thread blocks in the thread grid where the thread of dimension x is located, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and gridDim.x * blockDim.x represents the total number of threads in the thread grid where the thread of dimension x is located (if the number of elements in the vertex queue is much larger than the total number of threads (gridDim.x * blockDim.x), this grid-striding loop mode can enable each thread to process multiple data elements and make reasonable use of computing resources).

[0144] (5-8) Set the counter global_id4 to the global ID of the current thread obtained in step (5-3);

[0145] (5-9) Determine if global_id4 is less than the length of the back queue bque. If it is, get the 4th vertex bque[global_id4] in the back queue bque and then proceed to step (5-10). Otherwise, it means that all vertices in the back queue bque have been processed and then return to step (5-2).

[0146] (5-10) Determine whether the distance between the adjacent vertices in the adjacency list corresponding to the 4th global_id vertex bque[global_id4] in the backward queue bque is infinite. If so, update the distance between the adjacent vertices in the adjacency list corresponding to the 4th global_id vertex bque[global_id4] in the backward queue bque and the distance between the adjacent vertices in the backward distance array bdis, and add the adjacent vertices in the adjacency list corresponding to the 4th global_id vertex bque[global_id4] in the backward queue bque. Otherwise, it means that the 4th global_id vertex bque[global_id4] in the backward queue bque has been processed. Remove the vertex from the backward queue bque, and then proceed to step (5-11).

[0147] (5-11) Set global_id4 = global_id4 + gridDim.x * blockDim.x, and return to step (5-9);

[0148] (6) Copy the query results to CPU memory and output them.

[0149] In one embodiment, a computer device is provided, which may be a gateway. The computer device includes a processor, memory, a network interface, and a database connected via a system bus. The processor provides computing and control capabilities. The memory includes a non-volatile storage medium and internal memory. The non-volatile storage medium stores an operating system, computer programs, and a database. The internal memory provides an environment for the operation of the operating system and computer programs in the non-volatile storage medium. The database stores recorded IP addresses and corresponding MAC addresses of terminals in a local area network. The network interface communicates with external terminals via a network connection. When executed by the processor, the computer program implements a GPU-accelerated reachability query method.

[0150] In one embodiment, a computer device is provided, including a memory, a processor, and a computer program stored in the memory and executable on the processor, wherein the processor executes the computer program to implement all steps of the GPU-accelerated reachability query method of the present invention.

[0151] In one embodiment, a computer-readable storage medium is provided having a computer program stored thereon, which, when executed by a processor, implements all the steps of the GPU-accelerated reachability query method of the present invention.

[0152] Those skilled in the art will understand that all or part of the processes in the methods of the above embodiments can be implemented by a computer program instructing related hardware. The computer program can be stored in a non-volatile computer-readable storage medium, and when executed, it can include the processes of the embodiments of the above methods. Any references to memory, storage, databases, or other media used in the embodiments provided in this application can include non-volatile and / or volatile memory. Non-volatile memory can include read-only memory (ROM), programmable ROM (PROM), electrically programmable ROM (EPROM), electrically erasable programmable ROM (EEPROM), or flash memory. Volatile memory can include random access memory (RAM) or external cache memory. By way of illustration and not limitation, RAM is available in various forms, such as static RAM (SRAM), dynamic RAM (DRAM), synchronous DRAM (SDRAM), dual data rate SDRAM (DDRSDRAM), enhanced SDRAM (ESDRAM), synchronous link DRAM (SLDRAM), direct memory bus RAM (RDRAM), direct memory bus dynamic RAM (DRDRAM), and memory bus dynamic RAM (RDRAM), etc.

[0153] Those skilled in the art will readily understand that the above description is merely a preferred embodiment of the present invention and is not intended to limit the present invention. Any modifications, equivalent substitutions, and improvements made within the spirit and principles of the present invention should be included within the scope of protection of the present invention.< / int>

Claims

1. A GPU-accelerated reachability query method, characterized in that, Includes the following steps: (1) Receive graph data from the user, compress and preprocess the graph data to obtain compressed graph data; (2) Obtain the query points from the user, including the source vertex and the target vertex, configure the GPU environment, and copy the graph data and query points after compression in step (1) into the GPU memory in the configured GPU environment; (3) Calculate the average out-degree of the graph data copied to the GPU memory in step (2), and determine whether the average out-degree is greater than the preset threshold. If it is, proceed to step (4); otherwise, proceed to step (5). (4) Obtain the graph data copied to the GPU memory in step (2), and use the two-stage BFS algorithm to query whether the query point obtained in step (2) is reachable in the graph data as the query result, and then proceed to step (6); Step (4) specifically includes the following sub-steps: (4-1) Create a vertex queue vque, an edge queue eque, and a distance array dis. Initialize the vertex queue vque to the source vertex, initialize the edge queue eque to empty, and initialize the distances corresponding to the source vertex in the distance array dis to 0, and initialize the distances corresponding to the other vertices in the distance array dis to infinity. (4-2) Determine if the vertex queue vque is empty. If it is, take the source vertex and the target vertex as unreachable as the query result and the process ends. Otherwise, determine if the vertex queue vque contains the target vertex. If it is, take the source vertex and the target vertex as reachable as the query result and the process ends. Otherwise, proceed to step (4-3). (4-3) Get the global ID of the current thread; (4-4) Set the counter global_id1 to the global ID of the current thread obtained in step (4-3); (4-5) Determine if global_id1 is less than the length of the vertex queue vque. If it is, get the global_id1-th vertex vque[global-id1] in the vertex queue vque and then proceed to step (4-6). Otherwise, it means that all vertices in the vertex queue vque have been processed and then proceed to step (4-9). (4-6) Get the size of the adjacency list corresponding to the global_id1 vertex vque[global-id1] in the vertex queue vque. Based on the size of the adjacency list, get the prefix sum of the thread bundle corresponding to the thread of the vertex in the thread grid. Use this as the offset of the adjacent vertices of the global_id1 vertex vque[global-id1] in the corresponding adjacency list in the edge queue eque. Store all adjacent vertices in the edge queue eque. (4-7) Set global_id1 = global_id1 + gridDim.x * blockDim.x, and return to step (4-5); where gridDim.x represents the number of thread blocks in the thread grid where the thread of dimension x is located, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and gridDim.x * blockDim.x represents the total number of threads in the thread grid where the thread of dimension x is located; (4-8) Copy the edge queue eque obtained in step (4-6) to the shared memory, and use a hash table in the shared memory to perform collision detection processing on the edge queue eque to obtain the edge queue eque after collision detection processing. (4-9) Set the counter global_id2 to the global ID of the current thread obtained in step (4-3); (4-10) Determine if global_id2 is less than the length of the edge queue eque. If it is, get the global_id2-th vertex eque[global_id2] in the edge queue eque and then proceed to step (4-11). Otherwise, it means that all vertices in the edge queue eque have been processed and then return to step (4-2). (4-11) Determine whether the distance to the second global_id2 vertex eque[global_id2] in the distance array dis is infinite. If it is, update the distance to the second global_id2 vertex eque[global_id2] in the distance array dis and add the second global_id2 vertex eque[global_id2] to the vertex queue vque. Otherwise, proceed to step (4-12). (4-12) Set global_id2 = global_id2 + gridDim.x * blockDim.x, and return to step (4-10); (5) Obtain the graph data copied to the GPU memory in step (2), and use the bidirectional BFS algorithm to query whether the query point obtained in step (2) is reachable in the graph data as the query result, and then proceed to step (6). (6) Copy the query results to CPU memory and output them.

2. The GPU-accelerated reachability query method according to claim 1, characterized in that, Step (1) specifically includes the following sub-steps: (1-1) Receive graph data from the user, set counter i = 0, obtain the total number of edges and vertices in the graph data, and initialize array C and array R to empty; (1-2) Determine if i is equal to the total number of vertices in the graph data. If yes, use the obtained array C as the final adjacency list and array R as the final offset list. Array C and array R are the compressed graph data. The process ends. Otherwise, proceed to step (1-3). (1-3) Determine whether there is an adjacent vertex of the i-th vertex in the graph data. If so, add the adjacent vertices to array C in sequence, and record the position of the first adjacent vertex of the i-th vertex in the graph data in array C in array R. Then proceed to step (1-5). Otherwise, proceed to step (1-4). (1-4) Set the i-th element R[i] in array R to the (i-1)-th element R[i-1], and then proceed to step (1-5); (1-5) Set i = i + 1 and return to step (1-2).

3. The GPU-accelerated reachability query method according to claim 1 or 2, characterized in that, The average out-degree σ in step (3) is equal to: Where n represents the total number of vertices in the graph data, and out_deg(i) represents the out-degree of the i-th vertex in the graph data.

4. The GPU-accelerated reachability query method according to claim 3, characterized in that, Steps (4-6) specifically include the following sub-steps: (4-6-1) Get the size of the adjacency list corresponding to the global_id1 vertex vque[global-id1] in the vertex queue vque, set the counter i=0, initialize the array prefix to empty, which is used to record the prefix sum result of the thread corresponding to the global_id1 vertex in the thread grid, initialize the array nbr to empty, which is used to record the size of the adjacency list corresponding to the global_id1 vertex, and initialize the 0th element prefix[0] in the array prefix to 0; (4-6-2) Determine if i is equal to the size of array prefix - 1. If yes, use array prefix as the final prefix sum result table and the process ends. Otherwise, proceed to step (4-6-3). (4-6-3) Set the i-th element in the array prefix to prefix[i-1]+nbr[i-1], where nbr[i-1] represents the size of the adjacency list corresponding to the global_id1 vertex, and then proceed to step (4-6-4); (4-6-4) Set i = i + 1 and return to step (4-6-2).

5. The GPU-accelerated reachability query method according to claim 4, characterized in that, In step (4-3), the global ID of the thread is equal to: ID=blockIdx.x*blockDim.x+threadIdx.x Where blockIdx.x represents the number of the thread block corresponding to the thread of dimension x in the entire thread grid, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and threadIdx.x represents the number of the thread bundle in the thread grid in which the thread of dimension x is located.

6. The GPU-accelerated reachability query method according to claim 5, characterized in that, Step (5) specifically includes the following sub-steps: (5-1) Create a forward queue fque, a backward queue bque, a forward distance array fdis, and a backward distance array bdis. Initialize the forward queue fque as the source vertex, the backward queue bque as the target vertex, initialize the distance corresponding to the source vertex in the forward distance array fdis to 0, and initialize the distance corresponding to the other vertices to infinity. Initialize the distance corresponding to the target vertex in the backward distance array bdis to 0, and initialize the distance corresponding to the other vertices to infinity. (5-2) Determine whether there are any distances in the forward distance array fdis and the backward distance array bdis that are not infinite. If so, take the reachability of the source vertex and the target vertex as the query result and end the process. Otherwise, determine whether the forward queue fque and the backward queue bque are both empty. If so, take the unreachability of the source vertex and the target vertex as the query result and end the process. Otherwise, proceed to step (5-3). (5-3) Get the global ID of the current thread; (5-4) Set the counter global_id3 to the global ID of the current thread obtained in step (5-3); (5-5) Determine if global_id3 is less than the length of the forward queue fque. If it is, get the 3rd global_id3 vertex fque[global_id3] in the forward queue fque and then proceed to step (5-6). Otherwise, it means that all vertices in the forward queue fque have been processed and then proceed to step (5-8). (5-6) Determine whether the distance between the adjacent vertices in the adjacency list corresponding to the 3rd global_id vertex fque[global_id3] in the forward queue fque is infinite. If so, update the distance between the adjacent vertices in the adjacency list corresponding to the 3rd global_id vertex fque[global_id3] in the forward queue fque and the distance in the forward distance array fdis, and add the adjacent vertices in the adjacency list corresponding to the 3rd global_id vertex fque[global_id3] in the forward queue fque. Otherwise, it means that the 3rd global_id vertex fque[global_id3] in the forward queue fque has been processed. Remove the vertex from the forward queue fque and then proceed to step (5-7). (5-7) Set global_id3 = global_id3 + gridDim.x * blockDim.x, and return to step (5-5); where gridDim.x represents the number of thread blocks in the thread grid where the thread of dimension x is located, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and gridDim.x * blockDim.x represents the total number of threads in the thread grid where the thread of dimension x is located; (5-8) Set the counter global_id4 to the global ID of the current thread obtained in step (5-3); (5-9) Determine if global_id4 is less than the length of the back queue bque. If it is, get the 4th vertex bque[global_id4] in the back queue bque and then proceed to step (5-10). Otherwise, it means that all vertices in the back queue bque have been processed and then return to step (5-2). (5-10) Determine whether the distance between the adjacent vertices in the adjacency list corresponding to the 4th global_id vertex bque[global_id4] in the backward queue bque is infinite. If so, update the distance between the adjacent vertices in the adjacency list corresponding to the 4th global_id vertex bque[global_id4] in the backward queue bque and the distance between the adjacent vertices in the backward distance array bdis, and add the adjacent vertices in the adjacency list corresponding to the 4th global_id vertex bque[global_id4] in the backward queue bque. Otherwise, it means that the 4th global_id vertex bque[global_id4] in the backward queue bque has been processed. Remove the vertex from the backward queue bque, and then proceed to step (5-11). (5-11) Set global_id4 = global_id4 + gridDim.x * blockDim.x and return to step (5-9).

7. A GPU-accelerated reachability query system, characterized in that, include: The first module is used to receive graph data from the user, compress and preprocess the graph data to obtain compressed graph data. The second module is used to obtain query points from users, including source vertices and target vertices, configure the GPU environment, and copy the graph data compressed by the first module and the query points into the GPU memory in the configured GPU environment. The third module is used to calculate the average out-degree of the graph data copied to the GPU memory by the second module, and to determine whether the average out-degree is greater than a preset threshold. If it is, it proceeds to the fourth module; otherwise, it proceeds to the fifth module. The fourth module retrieves the graph data copied to the GPU memory by the second module, and uses a two-stage BFS algorithm to query whether the query point obtained by the second module is reachable from this graph data as the query result, which is then processed by the sixth module. The fourth module specifically executes the following sub-steps: (4-1) Create a vertex queue vque, an edge queue eque, and a distance array dis. Initialize the vertex queue vque to the source vertex, initialize the edge queue eque to empty, and initialize the distances corresponding to the source vertex in the distance array dis to 0, and initialize the distances corresponding to the other vertices in the distance array dis to infinity. (4-2) Determine if the vertex queue vque is empty. If it is, take the source vertex and the target vertex as unreachable as the query result and the process ends. Otherwise, determine if the vertex queue vque contains the target vertex. If it is, take the source vertex and the target vertex as reachable as the query result and the process ends. Otherwise, proceed to step (4-3). (4-3) Get the global ID of the current thread; (4-4) Set the counter global_id1 to the global ID of the current thread obtained in step (4-3); (4-5) Determine if global_id1 is less than the length of the vertex queue vque. If it is, get the global_id1-th vertex vque[global-id1] in the vertex queue vque and then proceed to step (4-6). Otherwise, it means that all vertices in the vertex queue vque have been processed and then proceed to step (4-9). (4-6) Get the size of the adjacency list corresponding to the global_id1 vertex vque[global-id1] in the vertex queue vque. Based on the size of the adjacency list, get the prefix sum of the thread bundle corresponding to the thread of the vertex in the thread grid. Use this as the offset of the adjacent vertices of the global_id1 vertex vque[global-id1] in the corresponding adjacency list in the edge queue eque. Store all adjacent vertices in the edge queue eque. (4-7) Set global_id1 = global_id1 + gridDim.x * blockDim.x, and return to step (4-5); where gridDim.x represents the number of thread blocks in the thread grid where the thread of dimension x is located, blockDim.x represents the number of threads contained in the thread block corresponding to the thread of dimension x, and gridDim.x * blockDim.x represents the total number of threads in the thread grid where the thread of dimension x is located; (4-8) Copy the edge queue eque obtained in step (4-6) to the shared memory, and use a hash table in the shared memory to perform collision detection processing on the edge queue eque to obtain the edge queue eque after collision detection processing. (4-9) Set the counter global_id2 to the global ID of the current thread obtained in step (4-3); (4-10) Determine if global_id2 is less than the length of the edge queue eque. If it is, get the global_id2-th vertex eque[global_id2] in the edge queue eque and then proceed to step (4-11). Otherwise, it means that all vertices in the edge queue eque have been processed and then return to step (4-2). (4-11) Determine whether the distance to the second global_id2 vertex eque[global_id2] in the distance array dis is infinite. If it is, update the distance to the second global_id2 vertex eque[global_id2] in the distance array dis and add the second global_id2 vertex eque[global_id2] to the vertex queue vque. Otherwise, proceed to step (4-12). The fifth module is used to obtain the graph data copied to the GPU memory by the second module, and use the bidirectional BFS algorithm to query whether the query point obtained by the second module is reachable in the graph data as the query result, and then transfer to the sixth module. The sixth module is used to copy the query results to CPU memory and output them.

Images