Local enhanced sparse matrix multiplied sparse vector calculation method and device

Abstract

The invention relates to the technical field of graph calculation and brain-like calculation, in particular to a locally enhanced sparse matrix multiplied sparse vector calculation method and device, and the method comprises the steps: dividing a column into a long column and a short column based on the number of non-zero elements in a sparse matrix column, and splitting the long column into a plurality of continuous long column fragments; reordering according to a target fragment list constructed by the short-column fragments and the long-column fragments, processing fragments in a new fragment list by utilizing each thread block on the GPU, caching the generated update pairs into a shared hash table of the target thread block, aggregating the update pairs meeting a preset condition in the shared hash table, and storing the aggregated update pairs in the shared hash table; and writing the aggregated update pairs back to the global output vector in batches to obtain a locally enhanced sparse matrix multiplied sparse vector calculation result. Therefore, the problems of insufficient data locality mining, increased memory access delay and reduction of the overall performance and resource utilization rate of the SpMSpV on the GPU due to relatively high overhead of a method based on sorting reduction in related technologies are solved.

Description

Technical Field

[0001] This application relates to the fields of graph computing and neuromorphic computing, and in particular to a method and apparatus for calculating sparse vector multiplication by locality enhancement sparse matrix. Background Technology

[0002] SpMSpV (Sparse Matrix-Sparse Vector Multiplication) is a fundamental operator in graph analysis and scientific computing. In graph computation, the adjacency matrix of a graph can be viewed as a sparse matrix A, and the set of currently active vertices can be represented as a sparse vector x. Algorithms such as BFS (Breadth-First Search), SSSP (Single-Source Shortest Path), PageRank, and connected component analysis are performed by iteratively calculating y = A * x or y = A ⊙ x with a mask. Some graph computation frameworks (such as implementations based on the GraphBLAS interface) uniformly express most graph algorithms as a series of SpMSpV and SpMV (Sparse Matrix-Vector Multiplication) operations.

[0003] Furthermore, in the field of neuromorphic computing, SpMSpV is also one of the core fundamental operators in the update process of SNN (Spiking Neural Network). Neuromorphic models typically employ an event-driven computational paradigm, where the firing of neurons (spikes) is highly sparsity in both time and space: at any given time step, only a small number of neurons generate spikes; the synaptic connection structure itself often exhibits a sparse topology. In this case, synaptic connections can be represented as sparse matrices, and the set of neurons generating spikes in the current time step can be represented as sparse vectors. The propagation of spikes in the network and the weighted accumulation of synapses are achieved through the multiplication of the sparse matrix and the sparse vector.

[0004] Specifically, the synaptic propagation process in SNNs can be formalized as v = W * s, where W is the sparse synaptic weight matrix, s is the sparse pulse vector at the current time step, and v is the accumulated synaptic current. This calculation is structurally highly consistent with column-driven SpMSpV, accessing only a small number of matrix columns triggered by pulses, avoiding traversal of the entire synaptic matrix, thus significantly reducing computational load and memory access overhead. In existing neuromorphic simulation and execution frameworks, SpMSpV has been widely adopted as a key underlying operator; for example, in BrainPy, SPMSPV is implemented using braininevent as a pulse propagation calculation method.

[0005] Secondly, SpMSpV is also widely used in the fields of artificial intelligence and scientific computing. For example, in the message passing phase of GNN (Graph Neural Network), it is necessary to aggregate features to neighbors based on the currently active vertex; multiplication between sparse matrices and sparse vectors also frequently occurs in sparse linear solutions, iterative methods, and preconditioner construction. Compared with traditional SpMV, SpMSpV can utilize the sparsity of the input vector, accessing only the matrix columns associated with non-zero vector elements, reducing invalid memory accesses and computations.

[0006] Existing SpMSpV implementations can be broadly categorized into row-driven (row-major, pull) and column-driven (col-major, push) operations. On GPUs (Graphics Processing Units), many graph computation and neuromorphic computing frameworks and research works employ column-driven SpMSpV as the core operator. For example, in GPU implementations based on the GraphBLAS interface, frameworks like Gunrock and GraphBLAST, message passing in BFS, PageRank, and GNN can all be decomposed into multiple SpMSpV operations. Current techniques distribute the non-zero elements of the input sparse vector to GPU threads or thread bundles, with each thread traversing one or more matrix column intervals and performing atomic addition on the output vector. To reduce conflicts, some works also write all (row index, partial sum) pairs into an intermediate array and then use sorting and merging to complete the aggregation.

[0007] However, the relevant technologies are mainly based on two frameworks: direct atomic write-back or write-sort-reduce. The sort-reduce method has a large overhead, resulting in insufficient data locality mining, increased memory access latency, and reduced overall performance and resource utilization of SpMSpV on GPUs, which urgently needs to be addressed. Summary of the Invention

[0008] This application is based on the inventor's understanding and insights into the following issues: SpMSpV is a fundamental operator in graph analysis and scientific computing. In graph computation, the adjacency matrix of a graph can be viewed as a sparse matrix A, and the set of currently active vertices can be represented as a sparse vector x. Algorithms such as BFS, SSSP, PageRank, and connected component analysis are performed by iteratively computing y = A * x or y = A ⊙ x with a mask. Some graph computation frameworks (such as implementations based on the GraphBLAS interface) uniformly express most graph algorithms as a series of SpMSpV and SpMV operations.

[0009] SpMSpV is also widely used in the fields of artificial intelligence and scientific computing. For example, in the message passing phase of graph neural networks (GNNs), features need to be aggregated to neighbors based on the currently active vertex; multiplication between sparse matrices and sparse vectors also frequently occurs in sparse linear solutions, iterative methods, and preconditioner construction. Compared with traditional SpMV, SpMSpV can take advantage of the sparsity of the input vector, accessing only the matrix columns associated with non-zero vector elements, reducing invalid memory accesses and computations.

[0010] In neuromorphic computing and SNNs, SpMSpV also constitutes a key fundamental computational paradigm. These models typically employ an event-driven execution approach. Neuronal spikes and synaptic connections exhibit high sparsity in both time and space. Synaptic weights can be represented as sparse matrices, and the current spike event can be represented as a sparse vector. The propagation and accumulation of spikes in the neural network are achieved through the multiplication of the sparse matrix and the sparse vector. Therefore, SpMSpV is widely used in neuromorphic simulation and execution frameworks to model synaptic propagation and neuronal input aggregation processes.

[0011] In terms of storage format, a sparse matrix A is typically represented using either CSR (Compressed Sparse Row) or CSC (Compressed Sparse Column): CSR format uses three arrays—row_ptr, indices, and values—to describe the row start and end positions, column indices, and values; CSC format uses col_ptr, indices, and values ​​to describe the column start and end positions, row indices, and values. A sparse vector x generally uses a compressed vector format, recording the index array idx of non-zero elements and the corresponding value array val, and may also include a bitmask if necessary.

[0012] like Figure 1 As shown, existing SpMSpV implementations can be broadly categorized into two types: row-driven (row-major, pull) and column-driven (col-major, push). The row-driven method is typically implemented in CSR format, iterating through each row r of the matrix in parallel, scanning all non-zero elements in that row (j = row_ptr[r]…row_ptr[r + 1]), reading the column index col = indices[j] for each non-zero element, checking if the input vector x is non-zero at that position, and if so, accumulating it into the row result res, finally writing back y[r]. This method does not require atomic updates to the output vector and is suitable for sequential or moderately parallel architectures such as CPUs (Central Processing Units).

[0013] Column-driven methods are typically implemented in CSC format. They iterate in parallel through each non-zero element x[i] of the input sparse vector x, obtaining the corresponding column col = idx[i] and vector value v_val = val[i]. Then, based on the row indices ind_list and matrix values ​​val_list within the range col_ptr[col]…col_ptr[col + 1], v_val is multiplied sequentially by the non-zero elements of that column, and the result is added to the corresponding position in the output vector y. Because each column element is "pushed" to multiple row positions, this method is also known as vector-driven or push-based SpMSpV.

[0014] On GPUs, many graph computing frameworks and research works use column-driven SpMSpV as the core operator. For example, in GPU implementations based on the GraphBLAS interface, frameworks such as Gunrock and GraphBLAST, message passing in BFS, PageRank, and GNN can all be decomposed into multiple SpMSpV operations. A common approach is to distribute the non-zero elements of the input sparse vector to GPU threads or thread bundles, with each thread traversing one or more matrix column intervals and performing atomic addition on the output vector. To reduce collisions, some works also write all (row index, partial sum) pairs into an intermediate array and then use sorting and merging to complete the aggregation.

[0015] For GPU architectures, existing SpMSpV implementations generally follow the two main approaches mentioned above: row-driven and column-driven, and further optimize load balancing and memory access. For example, some implementations allocate columns to different thread blocks according to column length, or use multi-level front queues to divide the workload; others circumvent the bottleneck of GPU atomic operations by sorting and segmenting update pairs (row, value). However, these methods still mainly revolve around the frameworks of "direct atomic write-back" or "write-sort-reduce," and there is still room for improvement in data locality mining and memory access latency hiding.

[0016] In light of the above background, existing SpMSpV implementations on GPUs generally have the following problems and shortcomings: (1) The memory access pattern is highly irregular. Whether it is row-driven or column-driven, it is necessary to access data in another array based on the index array, such as accessing x[col] or y[row]. This type of a[b[i]] index memory access makes the global memory access address distribution scattered, making it difficult to form merged access. The cache hit rate and L2 reuse rate are both low, which makes SpMSpV as a whole subject to memory bandwidth and memory access latency.

[0017] (2) The computation / memory access ratio is low, making it difficult to fully utilize GPU computing power. SpMSpV performs only one multiplication and one addition for each non-zero element, but requires at least two to three memory access operations (reading matrix indices and values, reading and writing vector elements), resulting in low overall arithmetic strength. Even with thousands of threads on the GPU, it is difficult to completely hide the long latency from global memory, causing a large number of cycles to be in a stall state.

[0018] (3) There are serious write conflicts in the output write-back. The typical implementation of column-driven SpMSpV on the GPU often directly performs atomicAdd on the output vector y. For graphs with power-law degree distribution or sparse matrices with uneven height, a few high out-degree vertices will generate a large number of update requests to write to the same row index. Frequent conflicts occur between thread bundles on the same address, which severely limits the throughput of atomic units.

[0019] (4) Sorting-based reduction methods have high overhead. To avoid a large number of atomic operations, some schemes write all intermediate results into a temporary array, then perform global sorting by row index, and then perform segmented reduction on the sorted results. This type of method can reduce atomic pressure when there are severe collisions, but it requires additional global memory space, reads and writes intermediate arrays multiple times, and introduces sorting overhead of O(nnz·log nnz) or O(nnz) level (depending on Radix Sort or Merge Sort), which is not economical for most matrices.

[0020] (5) Uneven column lengths lead to unbalanced load. Real graphs and scientific sparse matrices often have highly uneven column length distributions, with a few long columns containing a large number of non-zero elements, while many columns are very short. If the workload is simply distributed evenly by column or by the leading edge element, some thread blocks will need to process extremely long columns and run for a long time, while other thread blocks will finish early and be idle, resulting in low overall GPU resource utilization. Most existing load balancing strategies only consider column length or the number of non-zero elements, making it difficult to simultaneously address write-back pressure and local reuse.

[0021] (6) Insufficient locality mining. Existing methods often traverse columns according to the original storage order of the matrix or the order of the input vector, lacking specific optimization for row index locality. Even if shared memory is used to cache part of the results within the block, if the row indices corresponding to different columns are far apart and have low overlap, it is difficult to form high-density aggregation in the block-level cache or hash table, resulting in limited utilization value of shared memory.

[0022] (7) Insufficient utilization of GPU-specific memory access mechanisms. Some existing implementations still rely on a sequential mode of synchronous loading-computation-writeback, failing to fully utilize asynchronous copy instructions such as cp.async and double buffering techniques to hide memory access latency. They also lack a phased pipeline design oriented towards SpMSpV characteristics, resulting in a large number of cycles being wasted waiting for global memory.

[0023] In summary, existing GPU SpMSpV implementations either rely on a large number of global atomic operations or on global sorting and reduction, making it difficult to simultaneously solve the problems of poor memory access locality, heavy write-back conflicts, unbalanced column lengths, and high memory access latency. There is an urgent need for a new SpMSpV implementation method that is more suitable for GPU architecture, can enhance locality, and reduce write conflicts.

[0024] This application provides a method and apparatus for calculating locality-enhanced sparse matrix multiplication and sparse vector multiplication, in order to solve the problems in related technologies that are mainly based on two frameworks: direct atomic write-back or write-sort-reduce. The sort-reduce method has a large overhead, resulting in insufficient locality mining of data, increased memory access latency, and reduced overall performance and resource utilization of SpMSpV on GPUs.

[0025] The first aspect of this application provides a method for calculating sparse matrix multiplication and sparse vector multiplication with enhanced locality, comprising the following steps: dividing the target column into a target long column and a target short column based on the number of non-zero elements in the target column of the target sparse matrix, and further splitting the target long column into multiple consecutive long column segments; constructing a target segment list based on the target short column and the long column segments, and reordering the target segment list to obtain a new segment list; processing at least one segment in the new segment list using each thread block on the graphics processing unit (GPU) to generate at least one update pair, caching the at least one update pair in a shared hash table of the target thread block, aggregating at least one update pair in the shared hash table that meets a preset condition to obtain aggregated update pairs, and writing the aggregated update pairs back in batches to the global output vector to obtain the result of calculating sparse matrix multiplication and sparse vector multiplication with enhanced locality.

[0026] Optionally, in one embodiment of this application, the step of processing at least one fragment in the new fragment list using each thread block on the GPU to generate at least one update pair, caching the at least one update pair in a target shared hash table, and aggregating at least one update pair in the shared hash table that meets a preset condition to obtain an aggregated update pair includes: loading the index and value of the current fragment into a shared memory buffer according to the target asynchronous copy method, and reading the index and value of the current fragment from the shared memory buffer using a thread block to determine the current update pair using the index and value of the current fragment and an input sparse vector; caching the current update pair in the shared hash table of the corresponding thread block, aggregating the current update pair and other update pairs in the shared hash table, and when the aggregation process meets a preset condition, obtaining the index and value of the next fragment prefetched by each thread block when processing the current fragment; determining the next update pair using the index and value of the next fragment and the input sparse vector, caching the next update pair in the shared hash table of the corresponding thread block, and aggregating the next update pair, the current update pair, and other update pairs in the shared hash table to obtain the aggregated update pair.

[0027] Optionally, in one embodiment of this application, the step of writing the aggregated update pairs back to the global output vector in batches includes: in response to the shared hash table filling degree reaching a preset threshold, or in response to each thread block meeting a preset processing condition, sequentially traversing the shared hash table in each thread block using hash buckets to obtain the aggregated update pairs in the non-empty slots of the shared hash table; and writing the aggregated update pairs in the non-empty slots of the shared hash table back to the global output vector in batches.

[0028] Optionally, in one embodiment of this application, reordering the target fragment list to obtain a new fragment list includes: obtaining the row index of the first non-zero element of each fragment in the target fragment list; and using the row index of the first non-zero element of each fragment as a key for lightweight sorting to obtain the new fragment list.

[0029] Optionally, in one embodiment of this application, before caching the at least one update pair to the shared hash table of the target thread block, the method further includes: allocating a hash table of a preset size to each thread block in the shared memory to cache update pairs generated by at least one segment that the current thread block is responsible for; wherein, the hash table entry includes a row index key and an accumulated value, and the initial row index is set to -1 to indicate an empty slot.

[0030] A second aspect of this application provides a locality-enhanced sparse matrix-multiplication sparse vector computation apparatus, comprising: a splitting module, configured to divide the target column into a target long column and a target short column based on the number of non-zero elements in the target column of the target sparse matrix, and to split the target long column into multiple consecutive long column segments; an acquisition module, configured to construct a target segment list based on the target short column and the long column segments, and to reorder the target segment list to obtain a new segment list; and a processing module, configured to process at least one segment in the new segment list using each thread block on a graphics processing unit (GPU) to generate at least one update pair, cache the at least one update pair in a shared hash table of the target thread block, aggregate at least one update pair in the shared hash table that satisfies a preset condition to obtain aggregated update pairs, and write the aggregated update pairs back in batches to a global output vector to obtain the locality-enhanced sparse matrix-multiplication sparse vector computation result.

[0031] Optionally, in one embodiment of this application, the processing module includes: a first processing unit, configured to load the index and value of the current segment into a shared memory buffer according to the target asynchronous copy method, and read the index and value of the current segment from the shared memory buffer using a thread bundle, so as to determine the current update pair using the index and value of the current segment and the input sparse vector; a second processing unit, configured to cache the current update pair in the shared hash table of the corresponding thread block, aggregate the current update pair and other update pairs in the shared hash table, and when the aggregation process meets a preset condition, obtain the index and value of the next segment prefetched by each thread block when processing the current segment; and a first obtaining unit, configured to determine the next update pair using the index and value of the next segment and the input sparse vector, cache the next update pair in the shared hash table of the corresponding thread block, and aggregate the next update pair, the current update pair, and other update pairs in the shared hash table to obtain the aggregated update pair.

[0032] Optionally, in one embodiment of this application, the processing module includes: a second acquisition unit, configured to, in response to the shared hash table filling degree reaching a preset threshold, or in response to each thread block meeting a preset processing condition, sequentially traverse the shared hash table in each thread block using hash buckets to obtain the aggregated update pairs in the non-empty slots of the shared hash table; and a write-back unit, configured to batch write the aggregated update pairs in the non-empty slots of the shared hash table back to the global output vector.

[0033] Optionally, in one embodiment of this application, the acquisition module includes: a third acquisition unit, configured to acquire the row index of the first non-zero element of each segment in the target segment list; and a fourth acquisition unit, configured to perform lightweight sorting using the row index of the first non-zero element of each segment as a key to obtain the new segment list.

[0034] Optionally, in one embodiment of this application, the apparatus further includes: an allocation module, configured to allocate a hash table of a preset size to each thread block in shared memory to cache update pairs generated by at least one segment handled by the current thread block; wherein, the hash table entry includes a row index key and an accumulated value, and the initial row index is set to -1 to indicate an empty slot. A third aspect of this application provides an electronic device, including: a memory, a processor, and a computer program stored in the memory and executable on the processor, wherein the processor executes the program to implement the locality-enhanced sparse matrix multiplication sparse vector calculation method as described in the above embodiments.

[0035] A fourth aspect of this application provides a computer-readable storage medium storing a computer program that, when executed by a processor, implements the above-described method for calculating locality-enhanced sparse matrix multiplication and sparse vector multiplication.

[0036] A fifth aspect of this application provides a computer program product, including a computer program that, when executed, is used to implement the above-described method for calculating locality-enhanced sparse matrix multiplication and sparse vector multiplication.

[0037] This application embodiment can divide the sparse matrix column into long columns and short columns based on the number of non-zero elements in the column, and further split the long columns into multiple consecutive long column segments. The target segment list constructed based on the short and long column segments is reordered, and each thread block on the GPU processes the segments in the new segment list. The generated update pairs are cached in the shared hash table of the target thread block. Update pairs that meet preset conditions in the shared hash table are aggregated, and the aggregated update pairs are written back to the global output vector in batches to obtain sparse matrix multiplication of sparse vector calculation results with enhanced locality, effectively improving the overall performance and resource utilization of SpMSpV on the GPU. This solves the problem in related technologies where sorting-based reduction methods have high overhead, leading to insufficient data locality mining, increased memory access latency, and reduced overall performance and resource utilization of SpMSpV on the GPU.

[0038] Additional aspects and advantages of this application will be set forth in part in the description which follows, and in part will be obvious from the description, or may be learned by practice of this application. Attached Figure Description

[0039] The above and / or additional aspects and advantages of this application will become apparent and readily understood from the following description of the embodiments taken in conjunction with the accompanying drawings, wherein: Figure 1 This is a comparative diagram showing the basic calculation processes of row-driven sparse matrix multiplication and column-driven sparse matrix multiplication in related technologies. Figure 2 This is a flowchart of a method for calculating sparse vector multiplication using a locality-enhanced sparse matrix according to an embodiment of this application; Figure 3 A flowchart illustrating the locality enhancement sparse matrix multiplication by sparse vector in a specific embodiment of this application; Figure 4 This is a schematic diagram illustrating different load balancing strategies in a specific embodiment of this application; Figure 5 This is a schematic diagram illustrating the process of splitting a long column and rearranging segments according to a specific embodiment of this application; Figure 6 This is a schematic diagram of a locality-enhanced sparse matrix multiplication sparse vector computing device provided according to an embodiment of this application; Figure 7 This is a schematic diagram of the structure of an electronic device provided according to an embodiment of this application. Detailed Implementation

[0040] The embodiments of this application are described in detail below. Examples of these embodiments are shown in the accompanying drawings, wherein the same or similar reference numerals denote the same or similar elements or elements having the same or similar functions throughout. The embodiments described below with reference to the accompanying drawings are exemplary and intended to explain this application, and should not be construed as limiting this application.

[0041] The following describes a method and apparatus for calculating locality-enhanced sparse matrix multiplication and sparse vectors according to embodiments of this application, with reference to the accompanying drawings. Addressing the issues raised in the background section regarding related technologies primarily based on direct atomic write-back or write-sort-reduce frameworks, the sort-reduce method suffers from high overhead, leading to insufficient locality mining, increased memory access latency, and reduced overall performance and resource utilization of SpMSpV on GPUs. This application provides a method for calculating locality-enhanced sparse matrix multiplication and sparse vectors. In this method, columns are divided into long and short columns based on the number of non-zero elements in the sparse matrix, and the long columns are further divided into multiple consecutive long column segments. A target segment list constructed from the short and long column segments is reordered, and each thread block on the GPU processes the segments in the new segment list. The generated update pairs are cached in a shared hash table of the target thread block. Update pairs meeting preset conditions in the shared hash table are aggregated, and the aggregated update pairs are written back in batches to the global output vector to obtain the locality-enhanced sparse matrix multiplication and sparse vector calculation results, effectively improving the overall performance and resource utilization of SpMSpV on GPUs. This solves the problem that sorting-based reduction methods in related technologies have high overhead, leading to insufficient data locality mining, increased memory access latency, and reduced overall performance and resource utilization of SpMSpV on GPUs.

[0042] Specifically, Figure 2 This is a flowchart illustrating a method for calculating sparse vector multiplication using a locality-enhanced sparse matrix, as provided in an embodiment of this application.

[0043] like Figure 2 As shown, the method for calculating the locality-enhanced sparse matrix multiplication by sparse vector includes the following steps: In step S201, based on the number of non-zero elements in the target column of the target sparse matrix, the target column is divided into a target long column and a target short column, and the target long column is split into multiple consecutive long column segments.

[0044] In this embodiment, the target columns of the target sparse matrix are the columns of the sparse matrix A. First, the sparse matrix A can be stored in CSC format, with arrays col_ptr, indices, and values ​​recording the start and end positions, row indices, and values ​​of each column, respectively. The input sparse vector x is stored in a compressed format, recording the index array idx and the value array val for non-zero elements. The global output vector y is stored in a dense format.

[0045] It is understood that, in embodiments of this application, the target column can be divided into a target long column and a target short column based on the number of non-zero elements in the target column of the target sparse matrix, for example, as... Figure 3As shown, in this embodiment, the matrix columns can first be statistically analyzed and classified according to the number of non-zero elements in each column, and a long column threshold LEN_THRES is set, such as 128. When the number of non-zero elements in a column is greater than LEN_THRES, the column is considered long; otherwise, it is considered short. Short columns are processed at the original column granularity, while long columns are further divided into multiple continuous segments, and the number of non-zero elements in each segment does not exceed SPLIT_SIZE, such as 256.

[0046] Furthermore, embodiments of this application can convert the original column set of the sparse matrix into a list of segments consisting of "short column segments" and "long column segments". The splitting of long columns reduces the computation and write-back pressure on individual segments, which is beneficial for load balancing across different thread blocks and creates conditions for subsequent locality enhancement.

[0047] Therefore, compared with the prior art, the embodiments of this application can reduce the overhead of long columns while maintaining load balancing. Long column splitting divides a large number of non-zero elements that were originally concentrated in a single column into multiple segments and distributes them to different thread blocks for processing, avoiding the situation where a single block runs exclusively for a long time, thereby improving the overall GPU utilization. At the same time, a continuous index layout is still maintained within a single segment, without violating memory access locality.

[0048] In step S202, a target segment list is constructed based on the target short column and long column segments, and the target segment list is reordered to obtain a new segment list.

[0049] It is understood that embodiments of this application can construct a target fragment list based on target short and long fragments, and then, as follows: Figure 3 As shown, the obtained fragment list can be reordered in the following steps to obtain a new fragment processing order. The sorting only applies to fragment metadata, such as the fragment's column number, starting offset, and length, without moving the matrix body data, thus keeping preprocessing overhead low.

[0050] In this embodiment, after fragment rearrangement, logically adjacent fragments are closer in the row index space, significantly increasing the probability that different fragments update the same or adjacent rows. When the thread block traverses the sorted fragments and aggregates the results in shared memory in the following steps, the same or similar row indices are more likely to fall into the same local hash table, thereby improving the hash hit rate and local aggregation density.

[0051] For example, such as Figure 4 The diagram illustrates several typical load balancing strategies within the context of embodiments of this application, including direct allocation by column, coarse-grained allocation by the number of non-zero elements, and the fragment allocation method based on long column splitting adopted in this invention. Figure 4It is known that traditional strategies suffer from uneven load distribution in long column scenarios, while the embodiments of this application can achieve load balancing through column splitting.

[0052] Therefore, as Figure 5 As shown, embodiments of this application can split a long column containing a large number of non-zero elements into multiple consecutive segments, and sort these segments together with other short column segments according to the first row index to obtain a new processing order, thereby updating the spatial adjacency of segments in the same row or adjacent rows.

[0053] Optionally, in one embodiment of this application, reordering the target fragment list to obtain a new fragment list includes: obtaining the row index of the first non-zero element of each fragment in the target fragment list; and using the row index of the first non-zero element of each fragment as a key for lightweight sorting to obtain a new fragment list.

[0054] In actual implementation, the embodiments of this application can perform lightweight sorting on the obtained fragment list using the row index of the first non-zero element in the fragment as the key to obtain a new processing order for the fragment list, thereby keeping the preprocessing overhead low and improving the feasibility of locality-enhanced sparse matrix multiplication by sparse vector computation.

[0055] In step S203, at least one fragment in the new fragment list is processed by each thread block on the graphics processing unit (GPU) to generate at least one update pair. The at least one update pair is cached in the shared hash table of the target thread block. At least one update pair that meets the preset conditions in the shared hash table is aggregated to obtain aggregated update pairs. The aggregated update pairs are then written back to the global output vector in batches to obtain the locality-enhanced sparse matrix multiplication sparse vector calculation result.

[0056] In this embodiment, the GPU uses thread blocks as the basic unit of scheduling. Each thread block is responsible for processing several columns of segments and maintains a local hash table in shared memory to temporarily store updates to the output vector. Arrays for storing matrices and vectors, as well as necessary auxiliary arrays, are pre-allocated in global memory.

[0057] Specifically, such as Figure 3 As shown, in this embodiment of the application, the sorted list of fragments can be mapped to thread blocks on the GPU according to the fragment length and number, and a block mapping strategy is adopted: each thread block is responsible for a group of consecutive fragments, ensuring that the total number of non-zero numbers processed by each block is close, taking into account both load balancing and locality; the number of fragments each block is responsible for can also be adjusted according to the number of SMs and register / shared memory limits of different GPUs.

[0058] Furthermore, within each thread block, the work is further divided into segments at the thread bundle level. Threads within a thread bundle collaboratively traverse the indices and values ​​arrays within the segment, generate partial products, and call the local hash insertion function. By rationally designing the mapping relationship between threads and data, the merging degree of global memory access can be further improved.

[0059] It is understood that the embodiments of this application can utilize each thread block on the GPU to process one or more fragments in the new fragment list to generate one or more update pairs, and cache one or more update pairs in the shared hash table corresponding to the thread block. The update pairs in the shared hash table that meet the preset conditions are aggregated, for example, multiple update pairs targeting the same row index are merged in the block to obtain aggregated update pairs, and the aggregated update pairs are written back to the global output vector in batches to obtain the locality-enhanced sparse matrix multiplication sparse vector calculation results.

[0060] Therefore, the embodiments of this application can improve the overall performance and resource utilization of SpMSpV on GPUs without changing the global mathematical meaning of the sparse matrix. This is achieved through column classification and splitting, fragment rearrangement, block-level shared memory hash aggregation, and a data retrieval-computation-write-back pipeline, thereby reducing global atomic operations, increasing local aggregation density, and effectively hiding memory access latency.

[0061] It should be noted that the parameters such as hash table size, long column threshold LEN_THRES, fragment size SPLIT_SIZE, and FALLBACK_ITER in the embodiments of this application can be adjusted according to the GPU model, shared memory capacity, and typical matrix size to achieve a balance between hash aggregation efficiency, block occupancy rate, and preprocessing overhead.

[0062] In addition, the embodiments of this application can also construct a lightweight performance prediction model by statistically analyzing the distribution of non-zero numbers in each column and estimating the hash collision rate. When the predicted hash aggregation benefit is not obvious, it can degenerate into a traditional atomic or sorting implementation, thereby ensuring the robustness of the method on different matrices.

[0063] Optionally, in one embodiment of this application, each thread block on the GPU processes at least one fragment in a new fragment list to generate at least one update pair, and caches the at least one update pair in a target shared hash table. At least one update pair in the shared hash table that satisfies a preset condition is aggregated to obtain an aggregated update pair. This includes: loading the index and value of the current fragment into a shared memory buffer according to the target asynchronous copy method, and using a thread block to read the index and value of the current fragment from the shared memory buffer to determine the current update pair using the index and value of the current fragment and an input sparse vector; caching the current update pair in the shared hash table of the corresponding thread block, aggregating the current update pair and other update pairs in the shared hash table, and when the aggregation process satisfies a preset condition, obtaining the index and value of the next fragment prefetched by each thread block when processing the current fragment; determining the next update pair using the index and value of the next fragment and an input sparse vector, caching the next update pair in the shared hash table of the corresponding thread block, and aggregating the next update pair, the current update pair, and other update pairs in the shared hash table to obtain an aggregated update pair.

[0064] In some embodiments, this application can construct a three-stage pipeline of fetching, computation, and write-back within each thread block. Specifically, the thread block allocates double buffers bufA and bufB in shared memory for the indices and values ​​arrays, respectively, and utilizes asynchronous copy instructions such as cp.async provided by the GPU to prefetch data for the next segment while processing the current segment.

[0065] Specifically, when a thread block begins processing fragment k, its index and value are already loaded into a shared memory buffer in the background via asynchronous copying. During the computation phase, the thread block reads data from the buffer and performs hash insertion, aggregating the results into a shared memory hash table. When the aggregation is nearing completion, asynchronous loading of fragment k+1 is triggered in advance. Subsequently, when the hash table needs to be refreshed, a write-back phase is executed, writing the aggregation results back to global memory in batches. Through double buffering and asynchronous loading, the data transmission of another phase can overlap during the computation and write-back processes of one phase, effectively hiding global memory latency.

[0066] Therefore, compared with the prior art, the embodiments of this application can hide memory access latency by using a data fetch-computation-write-back pipeline. By building double buffering and asynchronous loading at the fragment level, the next batch of data can be loaded while performing hash operations and write-back during the SpMSpV calculation process, which significantly reduces the stall time caused by global memory access, so that the seemingly "computationally heavy" hash operation is mostly masked by memory access latency.

[0067] Optionally, in one embodiment of this application, writing the aggregated update pairs back to the global output vector in batches includes: in response to the shared hash table fill level reaching a preset threshold, or in response to each thread block meeting a preset processing condition, sequentially traversing the shared hash table in each thread block using hash buckets to obtain the aggregated update pairs in the non-empty slots of the shared hash table; and writing the aggregated update pairs in the non-empty slots of the shared hash table back to the global output vector in batches.

[0068] In actual implementation, such as Figure 3 As shown in the embodiment of this application, after the hash table fill level reaches a certain threshold or the thread block completes the calculation of the current batch of segments, the thread block traverses the shared hash table in each thread block in hash bucket order, and writes the aggregated update pairs (row_idx, value) results in non-empty slots back to the global output vector in batches. Since the entries in the hash table are accessed in bucket index order, the distance between adjacent write-back addresses is relatively continuous, which is beneficial to improving the merging degree of global memory write-back. After the write-back is completed, the corresponding slot is reset to empty for use by subsequent segments.

[0069] Therefore, compared with existing technologies, the embodiments of this application can significantly reduce global atomic operations and improve bandwidth utilization. By using a shared memory hash table for local aggregation within the thread block, multiple updates targeting the same row index are merged within the block before being written back to global memory, effectively reducing the dependence on global atomic operations. On typical graphs and sparse matrices, compared with column-driven SpMSpV with direct atomic write-back, the pressure on atomic units and the number of global write-backs are significantly reduced, improving bandwidth utilization.

[0070] Optionally, in one embodiment of this application, before caching at least one update pair to the shared hash table of the target thread block, the method further includes: allocating a hash table of a preset size to each thread block in the shared memory to cache update pairs generated by at least one segment that the current thread block is responsible for; wherein, the hash table entry includes a row index key and an accumulated value, and the initial row index is set to -1 to indicate an empty slot.

[0071] As one possible implementation, each thread block in this embodiment allocates a fixed-size hash table in shared memory to cache the update pairs (row_idx, value) generated by the current block's responsible fragment. Each hash table entry includes two fields: a row index key and an accumulated value. The initial row index is set to -1 to indicate an empty slot.

[0072] Specifically, hash insertion uses a linear probing method. The thread calculates the initial hash position based on the row index and uses atomicCAS to try to occupy an empty slot or find a slot with the same row index. If successful, it adds part of the value to the value field of the slot through atomic addition. If the number of probes exceeds the preset FALLBACK_ITER threshold, the hash table is considered to be overly conflicted, and the current update degenerates into directly performing atomicAdd on the global output vector to ensure the correctness of the algorithm.

[0073] For example, the embodiments of this application can achieve stable performance improvements on a variety of real-world datasets. On an NVIDIA A100 GPU, tests were conducted on over 100 large-scale Web graphs from Konect / LAW and over 200 scientific computing matrices from SuiteSparse, all of which have at least 5 million non-zero elements. Under four input vector sparsity levels (0.01, 0.05, 0.10, 0.20), compared to seven baseline methods—cuSPARSE, two representative row-driven SpMSpV methods, and four representative column-driven SpMSpV methods (including atomic write-back and sorted write-back)—the embodiments of this application achieve a geometric mean speedup of approximately 1.41x on Web graphs, with a maximum of approximately 3.42x, relative to the best "one-in-seven" baseline; and a geometric mean speedup of approximately 1.13x on SuiteSparse matrices, with a maximum of approximately 2.55x. The embodiments of this application achieve the best performance among all implementations in approximately 90% of the matrices on the Web graph, and this proportion is approximately 60% on SuiteSparse. Therefore, the embodiments of this application are not only effective on a small number of specific matrices, but also have a stable advantage over a wide range of real-world workloads.

[0074] Therefore, compared with existing technologies, the embodiments of this application can enhance row index locality and increase aggregation density. By splitting long columns and rearranging all segments according to row index order, the same or similar row indexes are more likely to appear in clusters in time and space, making the aggregation in the hash table "dense," and the same slot or adjacent slots are reused. Experimental observations show that the local overlap rate and write-back merge factor of this application are significantly better than the baseline implementation without rearrangement on typical Web graphs.

[0075] Furthermore, the embodiments of this application have a wide range of implementations and are easy to integrate. This application is compatible with conventional CSC format and sparse vector format in terms of interface, does not depend on specific matrix patterns or graph structures, and can be directly embedded into graph computing frameworks or sparse linear algebra libraries, providing a unified acceleration operator for various applications such as BFS, PageRank, and GNN message passing. It can also be implemented as an efficient neuromorphic operator library. At the same time, through parameter tuning and performance prediction models, this application has good adaptability to sparse matrices of different sizes and types.

[0076] The locality-enhanced sparse matrix multiplication and sparse vector calculation method proposed in this application can divide the sparse matrix column into long columns and short columns based on the number of non-zero elements in the column, and further split the long columns into multiple consecutive long column segments. The target segment list constructed based on the short and long column segments is reordered, and each thread block on the GPU processes the segments in the new segment list. The generated update pairs are cached in the shared hash table of the target thread block. Update pairs that meet preset conditions in the shared hash table are aggregated, and the aggregated update pairs are written back to the global output vector in batches to obtain the locality-enhanced sparse matrix multiplication and sparse vector calculation results, effectively improving the overall performance and resource utilization of SpMSpV on the GPU. This solves the problem that the sorting-reduction-based methods in related technologies have high overhead, leading to insufficient data locality mining, increased memory access latency, and reduced overall performance and resource utilization of SpMSpV on the GPU.

[0077] Next, with reference to the accompanying drawings, a locality-enhanced sparse matrix multiplication sparse vector computing apparatus according to an embodiment of this application is described.

[0078] Figure 6 This is a block diagram of a locality-enhanced sparse matrix multiplication sparse vector computation device according to an embodiment of this application.

[0079] like Figure 6 As shown, the locality-enhanced sparse matrix multiplication sparse vector computing device 10 includes: a splitting module 100, an acquisition module 200, and a processing module 300.

[0080] Specifically, the splitting module 100 is used to divide the target column into a target long column and a target short column based on the number of non-zero elements in the target column of the target sparse matrix, and to split the target long column into multiple consecutive long column segments.

[0081] The acquisition module 200 is used to construct a target fragment list based on the target short column and long column fragments, and to reorder the target fragment list to obtain a new fragment list.

[0082] The processing module 300 is used to process at least one fragment in a new fragment list using each thread block on the graphics processing unit (GPU) to generate at least one update pair, cache at least one update pair in a shared hash table of the target thread block, aggregate at least one update pair in the shared hash table that meets a preset condition to obtain aggregated update pairs, and write the aggregated update pairs back to the global output vector in batches to obtain the locality-enhanced sparse matrix multiplication sparse vector calculation result.

[0083] Optionally, in one embodiment of this application, the processing module 300 includes: a first processing unit, a second processing unit, and a first acquisition unit.

[0084] The first processing unit is used to load the index and value of the current segment into a shared memory buffer according to the asynchronous copy method of the target, and use a thread bundle to read the index and value of the current segment from the shared memory buffer, so as to determine the current update pair using the index and value of the current segment and the input sparse vector; The second processing unit is used to cache the current update pair in the shared hash table of the corresponding thread block, aggregate the current update pair and other update pairs in the shared hash table, and when the aggregation process meets the preset conditions, obtain the index and value of the next segment prefetched by each thread block when processing the current segment; The first acquisition unit is used to determine the next update pair by using the index and value of the next segment and the input sparse vector, and cache the next update pair in the shared hash table of the corresponding thread block, and aggregate the next update pair, the current update pair and other update pairs in the shared hash table to obtain the aggregated update pair.

[0085] Optionally, in one embodiment of this application, the processing module 300 includes a second acquisition unit and a write-back unit.

[0086] The second acquisition unit is used to sequentially traverse the shared hash table in each thread block using hash buckets in response to the shared hash table filling degree reaching a preset threshold, or in response to each thread block meeting a preset processing condition, in order to obtain the aggregated update pairs in the non-empty slots of the shared hash table.

[0087] The write-back unit is used to batch write the aggregated update pairs from the non-empty slots of the shared hash table back to the global output vector.

[0088] Optionally, in one embodiment of this application, the acquisition module 200 includes a third acquisition unit and a fourth acquisition unit.

[0089] The third acquisition unit is used to acquire the row index of the first non-zero element of each segment in the target segment list.

[0090] The fourth acquisition unit is used to perform lightweight sorting using the row index of the first non-zero element of each fragment as the key, in order to obtain a new list of fragments.

[0091] Optionally, in one embodiment of this application, the apparatus 10 of this application embodiment further includes: an allocation module.

[0092] The allocation module is used to allocate a hash table of a preset size to each thread block in the shared memory to cache update pairs generated by at least one segment that the current thread block is responsible for; wherein, the hash table entries include row index keys and accumulated values, and the initial row index is set to -1 to indicate an empty slot.

[0093] It should be noted that the foregoing explanation of the embodiment of the locality-enhanced sparse matrix multiplication sparse vector calculation method also applies to the locality-enhanced sparse matrix multiplication sparse vector calculation device of this embodiment, and will not be repeated here.

[0094] The locality-enhanced sparse matrix multiplication sparse vector computation device proposed in this application can divide the columns of a sparse matrix into long columns and short columns based on the number of non-zero elements in the columns, and further split the long columns into multiple consecutive long column segments. The target segment list constructed based on the short and long column segments is reordered, and each thread block on the GPU processes the segments in the new segment list. The generated update pairs are cached in the shared hash table of the target thread block. Update pairs that meet preset conditions in the shared hash table are aggregated, and the aggregated update pairs are written back in batches to the global output vector to obtain the locality-enhanced sparse matrix multiplication sparse vector computation result, effectively improving the overall performance and resource utilization of SpMSpV on the GPU. This solves the problem in related technologies where sorting-based reduction methods have high overhead, leading to insufficient data locality mining, increased memory access latency, and reduced overall performance and resource utilization of SpMSpV on the GPU.

[0095] Figure 7 A schematic diagram of the structure of an electronic device provided in an embodiment of this application. The electronic device may include: The memory 701, the processor 702, and the computer program stored on the memory 701 and executable on the processor 702.

[0096] When the processor 702 executes the program, it implements the locality-enhanced sparse matrix multiplication sparse vector calculation method provided in the above embodiments.

[0097] Furthermore, electronic devices also include: Communication interface 703 is used for communication between memory 701 and processor 702.

[0098] The memory 701 is used to store computer programs that can run on the processor 702.

[0099] The memory 701 may include high-speed RAM memory, and may also include non-volatile memory, such as at least one disk storage device.

[0100] If the memory 701, processor 702, and communication interface 703 are implemented independently, then the communication interface 703, memory 701, and processor 702 can be interconnected via a bus to complete communication between them. The bus can be an Industry Standard Architecture (ISA) bus, a Peripheral Component Interconnect (PCI) bus, or an Extended Industry Standard Architecture (EISA) bus, etc. Buses can be categorized into address buses, data buses, control buses, etc. For ease of representation, Figure 7 The bus is represented by a single thick line, but this does not mean that there is only one bus or one type of bus.

[0101] Optionally, in a specific implementation, if the memory 701, processor 702, and communication interface 703 are integrated on a single chip, then the memory 701, processor 702, and communication interface 703 can communicate with each other through an internal interface.

[0102] The processor 702 may be a central processing unit (CPU), an application specific integrated circuit (ASIC), or one or more integrated circuits configured to implement the embodiments of this application.

[0103] This embodiment also provides a computer-readable storage medium storing a computer program that, when executed by a processor, implements the above-described method for calculating sparse matrix multiplication of locality-enhanced sparse vectors.

[0104] This embodiment also provides a computer program product, including a computer program that, when executed, is used to implement the above-described method for calculating sparse matrix multiplication and sparse vector multiplication with enhanced locality.

[0105] In the description of this specification, the references to terms such as "one embodiment," "some embodiments," "example," "specific example," or "some examples," etc., indicate that a specific feature, structure, material, or characteristic described in connection with that embodiment or example is included in at least one embodiment or example of this application. In this specification, the illustrative expressions of the above terms do not necessarily refer to the same embodiment or example. Furthermore, the specific features, structures, materials, or characteristics described may be combined in any suitable manner in one or more embodiments or examples. Moreover, without contradiction, those skilled in the art can combine and integrate the different embodiments or examples described in this specification, as well as the features of different embodiments or examples.

[0106] Furthermore, the terms "first" and "second" are used for descriptive purposes only and should not be construed as indicating or implying relative importance or implicitly specifying the number of technical features indicated. Thus, a feature defined as "first" or "second" may explicitly or implicitly include at least one of that feature. In the description of this application, "N" means at least two, such as two, three, etc., unless otherwise explicitly specified.

[0107] Any process or method described in the flowchart or otherwise herein can be understood as representing a module, segment, or portion of code comprising one or N executable instructions for implementing custom logic functions or processes, and the scope of the preferred embodiments of this application includes additional implementations in which functions may be performed not in the order shown or discussed, including substantially simultaneously or in reverse order depending on the functions involved, as should be understood by those skilled in the art to which embodiments of this application pertain.

[0108] The logic and / or steps represented in the flowchart or otherwise described herein, for example, can be considered as a sequenced list of executable instructions for implementing logical functions, and can be embodied in any computer-readable medium for use by, or in conjunction with, an instruction execution system, apparatus, or device (such as a computer-based system, a processor-included system, or other system that can fetch and execute instructions from, an instruction execution system, apparatus, or device). For the purposes of this specification, "computer-readable medium" can be any means that can contain, store, communicate, propagate, or transmit programs for use by, or in conjunction with, an instruction execution system, apparatus, or device. More specific examples (a non-exhaustive list) of computer-readable media include: an electrical connection having one or more wires (electronic device), a portable computer disk drive (magnetic device), random access memory (RAM), read-only memory (ROM), erasable and editable read-only memory (EPROM or flash memory), fiber optic devices, and portable optical disc read-only memory (CDROM). Alternatively, the computer-readable medium may be paper or other suitable media on which the program can be printed, since the program can be obtained electronically by optically scanning the paper or other medium, followed by editing, interpreting, or otherwise processing as necessary, and then stored in a computer memory.

[0109] It should be understood that the various parts of this application can be implemented using hardware, software, firmware, or a combination thereof. In the above embodiments, the N steps or methods can be implemented using software or firmware stored in memory and executed by a suitable instruction execution system. For example, if implemented in hardware, as in another embodiment, it can be implemented using any one or a combination of the following techniques known in the art: discrete logic circuits having logic gates for implementing logical functions on data signals, application-specific integrated circuits (ASICs) having suitable combinational logic gates, programmable gate arrays (PGAs), field-programmable gate arrays (FPGAs), etc.

[0110] Those skilled in the art will understand that all or part of the steps of the methods described in the above embodiments can be implemented by a program instructing related hardware. The program can be stored in a computer-readable storage medium, and when executed, it includes one or a combination of the steps of the method embodiments.

[0111] Furthermore, the functional units in the various embodiments of this application can be integrated into a processing module, or each unit can exist physically separately, or two or more units can be integrated into a module. The integrated module can be implemented in hardware or as a software functional module. If the integrated module is implemented as a software functional module and sold or used as an independent product, it can also be stored in a computer-readable storage medium.

[0112] The storage medium mentioned above can be a read-only memory, a disk, or an optical disk, etc. Although embodiments of this application have been shown and described above, it is understood that the above embodiments are exemplary and should not be construed as limiting this application. Those skilled in the art can make changes, modifications, substitutions, and variations to the above embodiments within the scope of this application.

Claims

1. A method for calculating sparse vector multiplication using a locality-enhanced sparse matrix, characterized in that, Includes the following steps: Based on the number of non-zero elements in the target column of the target sparse matrix, the target column is divided into a target long column and a target short column, and the target long column is further split into multiple consecutive long column segments; A target segment list is constructed based on the target short column and the target long column segments, and the target segment list is reordered to obtain a new segment list; At least one fragment in the new fragment list is processed by each thread block on the graphics processing unit (GPU) to generate at least one update pair. The at least one update pair is cached in the shared hash table of the target thread block. At least one update pair in the shared hash table that meets a preset condition is aggregated to obtain aggregated update pairs. The aggregated update pairs are then written back to the global output vector in batches to obtain the locality-enhanced sparse matrix multiplication sparse vector calculation result.

2. The method of claim 1, wherein, The process of using each thread block on the GPU to process at least one fragment in the new fragment list to generate at least one update pair, caching the at least one update pair in the target's shared hash table, and aggregating at least one update pair in the shared hash table that meets preset conditions to obtain aggregated update pairs includes: According to the target asynchronous copy method, the index and value of the current segment are loaded into the shared memory buffer, and the thread bundle is used to read the index and value of the current segment from the shared memory buffer, so as to determine the current update pair using the index and value of the current segment and the input sparse vector; The current update pair is cached in the shared hash table of the corresponding thread block. The current update pair and other update pairs in the shared hash table are aggregated. When the aggregation process meets the preset conditions, the index and value of the next segment prefetched by each thread block when processing the current segment are obtained. The next update pair is determined by using the index and value of the next segment with the input sparse vector, and the next update pair is cached in the shared hash table of the corresponding thread block. The next update pair, the current update pair, and other update pairs in the shared hash table are aggregated to obtain the aggregated update pair.

3. The method of claim 1, wherein, The step of writing the aggregated update pairs back to the global output vector in batches includes: In response to the shared hash table filling degree reaching a preset threshold, or in response to each thread block meeting a preset processing condition, the shared hash table in each thread block is sequentially traversed using hash buckets to obtain the aggregated update pairs in the non-empty slots of the shared hash table; The aggregated update pairs in the non-empty slots of the shared hash table are batched and written back to the global output vector.

4. The method of claim 1, wherein, The step of reordering the target fragment list to obtain a new fragment list includes: Obtain the row index of the first non-zero element of each segment in the target segment list; Lightly sort the new list of segments by using the row index of the first non-zero element of each segment as the key.

5. The method of claim 1, wherein, Before caching the at least one update pair in the shared hash table of the target thread block, the method further includes: Allocate a hash table of a preset size for each thread block in shared memory to cache update pairs generated by at least one segment that the current thread block is responsible for; The hash table entry includes a row index key and an accumulated value, with the initial row index set to -1 to indicate an empty slot.

6. A locality-enhanced sparse matrix multiplication and sparse vector computation device, characterized in that, include: The splitting module is used to divide the target column into a target long column and a target short column based on the number of non-zero elements in the target column of the target sparse matrix, and to split the target long column into multiple consecutive long column segments; The acquisition module is used to construct a target segment list based on the target short column and the target long column segments, and to reorder the target segment list to obtain a new segment list; The processing module is configured to process at least one fragment in the new fragment list using each thread block on the graphics processing unit (GPU) to generate at least one update pair, cache the at least one update pair in the shared hash table of the target thread block, aggregate at least one update pair in the shared hash table that meets a preset condition to obtain aggregated update pairs, and write the aggregated update pairs back to the global output vector in batches to obtain the locality-enhanced sparse matrix multiplication sparse vector calculation result.

7. The apparatus of claim 6, wherein, The processing module includes: The first processing unit is configured to load the index and value of the current segment into a shared memory buffer according to the asynchronous copy method of the target, and use a thread bundle to read the index and value of the current segment from the shared memory buffer, so as to determine the current update pair using the index and value of the current segment and the input sparse vector; The second processing unit is used to cache the current update pair in the shared hash table of the corresponding thread block, aggregate the current update pair and other update pairs in the shared hash table, and when the aggregation process meets the preset conditions, obtain the index and value of the next segment prefetched by each thread block when processing the current segment; The acquisition unit is configured to determine the next update pair using the index and value of the next segment and the input sparse vector, cache the next update pair in the shared hash table of the corresponding thread block, and aggregate the next update pair, the current update pair, and other update pairs in the shared hash table to obtain the aggregated update pair.

8. An electronic device, comprising: include: The memory, the processor, and the computer program stored in the memory and capable of running on the processor, wherein the processor executes the program to implement the locality-enhanced sparse matrix multiplication sparse vector computation method as described in any one of claims 1-5.

9. A computer-readable storage medium having a computer program stored thereon, characterized in that, The program is executed by the processor to implement the locality-enhanced sparse matrix multiplication sparse vector computation method as described in any one of claims 1-5.

10. A computer program product, comprising a computer program, characterized in that, The computer program is executed by a processor to implement the locality-enhanced sparse matrix multiplication sparse vector computation method as described in any one of claims 1-5.

Images