AsyncSparse: Accelerating Sparse Matrix-Matrix Multiplication on Asynchronous GPU Architectures

Abstract: Sparse Matrix-Matrix Multiplication (SpMM) is a fundamental kernel across scientific computing and machine learning. While prior work accelerates SpMM using Tensor Cores, no existing sparse kernel exploits the asynchronous features of modern GPU architectures, such as NVIDIA's Tensor Memory Accelerator (TMA) and warp specialization. This work systematically studies how these features impact SpMM performance and introduces two co-designed kernels. For structured sparsity, we optimize a warp-specialized producer-consumer pipeline overlapping TMA data transfer with WGMMA computation using Block Compressed Sparse Row (BCSR) format. For irregular sparsity, we design a Window Compressed Sparse Row (WCSR) kernel that loads the sparse operand via TMA and splits large row-windows across thread blocks for load balancing. Our WCSR kernel outperforms all prior SpMM kernels on SuiteSparse matrices (1.47x over AccSpMM, 6.24x over cuSPARSE). Our BCSR kernel achieves a combined 2.66x end-to-end speedup on Qwen2.5-7B prefill at 90% block sparsity with 64K tokens over cuDNN/cuBLAS.

• The paper presents asynchronous GPU kernel designs for sparse matrix multiplication using BCSR and WCSR formats, achieving significant speedups.
• It details architectural innovations like TMA, WGMMA, and warp specialization to optimize memory transfer and load balance across pipelines.
• Empirical evaluations show up to 4.41× speedup over CUDA-core baselines and notable improvements in large language model inference performance.

AsyncSparse: Accelerating Sparse Matrix-Matrix Multiplication on Asynchronous GPU Architectures

Architectural Foundations and Sparse Format Design

The paper targets optimizing Sparse Matrix-Matrix Multiplication (SpMM)—a foundational computation in scientific computing, graph analytics, and machine learning. The work leverages hardware specialization in the NVIDIA Hopper architecture, centering on three architectural innovations: Tensor Memory Accelerator (TMA), Warpgroup Matrix-Multiply-Accumulate (WGMMA), and warp specialization for producer-consumer pipelines. Figure 1 provides a schematic overview, illustrating the Hopper architecture's organization of SMs, TMA units, Tensor Cores, shared memory, and the high-bandwidth memory subsystem.

Figure 1: NVIDIA Hopper GPU architecture, highlighting TMA units, Tensor Core blocks, shared memory, L2 cache, and HBM3.

Two primary sparse storage formats facilitate hardware exploitation:

• Block Compressed Sparse Row (BCSR): Tiles the sparse matrix into fixed-size blocks, storing only non-zero blocks, enabling contiguous memory access patterns amenable to TMA bulk transfer.
• Window Compressed Sparse Row (WCSR): Groups rows into windows, compresses non-zero columns per window, reducing zero-padding overhead in highly irregular sparse matrices.

  Figure 2: Schematic comparison of BCSR and WCSR formats, illustrating block and window granularity, contiguous vs. indirect memory access, and padding trade-offs.

Asynchronous Execution Paradigm: Kernel Design and Optimization

The AsyncSparse kernels are engineered to maximize utilization of Hopper's asynchronous features. The workflow is decomposed into specialized kernel designs for BCSR and WCSR formats:

• BCSR Warp-Specialized Pipeline: Employs a producer warpgroup to issue TMA loads; consumer warpgroups execute WGMMA instructions over tiles in a multi-stage circular buffer. The producer-consumer pipeline leverages asynchronous data movement and compute, eliminating idle Tensor Core phases.

  Figure 3: Pipeline architecture for BCSR kernel, detailing warpgroup specialization, TMA-driven shared memory buffer, and concurrent WGMMA execution.

• WCSR Load-Balanced Kernel: Uses cooperative threads for indirect dense operand gathering (necessitated by non-contiguous column access), with task-based decomposition to split windows for improved load balance across thread blocks.

The design incorporates explicit synchronization via Hopper’s mbarrier primitives and dynamic register allocation through PTX instructions. An ablation study demonstrates that incremental adoption of WGMMA, TMA, and warp specialization generates up to a 4.41× speedup over the CUDA-core baseline, accounting for over 98% of total performance improvement.

Figure 4: Execution timeline comparison between synchronous cooperative loads and asynchronous TMA pipelines, elucidating hardware concurrency.

Empirical Performance Evaluation

Extensive benchmarking on 414 SuiteSparse matrices showcases AsyncSparse’s throughput and scalability. The WCSR kernel demonstrates geometric-mean speedups of 1.47× over AccSpMM, 6.24× over cuSPARSE at dense matrix width N=1024. The box plot distribution confirms consistently higher median speedups for matrices exceeding 0.5% density.

Figure 5: Distribution of speedups vs. cuSPARSE for N=256, 512, 1024, with clear dominance of WCSR and BCSR kernels.

A detailed breakdown isolates each architectural feature’s contribution; WGMMA, TMA, and warp specialization yield the majority of gains, while persistent kernel launches and TMA multicast, beneficial for dense GEMM, regress in sparse SpMM due to co-scheduling constraints and increased synchronization overhead.

Figure 6: Performance breakdown from the ablation study, showing cumulative speedup as optimizations are enabled, normalized to cuSPARSE.

Tile sizing analysis reveals performance sensitivity to the tile width parameter ($\mathit{WGMMA\_N}$), with optimal throughput achieved when the tile size divides the dense matrix dimension evenly, minimizing padding overhead.

Figure 7: Throughput variance in the warp-specialized kernel as a function of $\mathit{WGMMA\_N}$, showing peak performance at dimension-aligned tile sizes.

LLM Integration and End-to-End Latency Reduction

The AsyncSparse BCSR kernel was integrated into a Qwen2.5-7B transformer model to benchmark prefill latency reduction at varying sparsity levels and sequence lengths. Combined with sparse attention acceleration (MInference), the approach realizes up to 2.66× end-to-end prefill speedup at a 64K-token context with 90% FFN block sparsity.

Figure 8: End-to-end prefill speedup in Qwen2.5-7B, juxtaposing sparse attention (MInference), BCSR FFN speedup, and combined performance at 64K tokens.

Theoretical and Practical Implications

AsyncSparse provides several key insights:

• Leveraging asynchronous execution primitives, notably TMA and WGMMA, produces substantial speedups for SpMM workloads; architectural co-design is critical for achieving memory/computation overlap.
• BCSR and WCSR formats exhibit complementary strengths: BCSR is competitive at higher matrix densities and aligns with hardware tile semantics, while WCSR minimizes zero-padding for highly irregular sparsity.
• Persistent kernel and multicast techniques—effective for dense linear algebra—do not directly translate to sparse operations, primarily due to load imbalance and reduced scheduler flexibility.
• Integration into LLM prefill chains demonstrates practical impact, with empirical speedup scaling favorably with increasing sparsity and sequence length.

On a theoretical plane, the results imply asynchronous hardware features will become increasingly central to sparse computation optimization as neural workloads scale and deployments grow more memory-bound. Future work may extend toward finer granularity formats, improved preprocessing/reordering for higher fill-ratios, FP8 support, and multi-GPU scaling.

Conclusion

AsyncSparse establishes a high-performance paradigm for SpMM on NVIDIA Hopper, validating the importance of hardware/software co-design in sparse computation. The suite of kernels achieves state-of-the-art throughput on diverse benchmark matrices and demonstrates substantial acceleration in LLM inference via block sparsity and asynchronous execution. The work highlights the nuanced interplay between format design, hardware specialization, and scheduling—pointing toward promising directions in future sparse computing research and practical deployment (2604.17834).