ClusterFusion: Expanding Operator Fusion Scope for LLM Inference via Cluster-Level Collective Primitive
Authors: Xinhao Luo
Zihan Liu
Yangjie Zhou
Shihan Fang
Ziyu Huang
Yu Feng
Chen Zhang
Shixuan Sun
Zhenzhe Zheng
Jingwen Leng
Minyi Guo
Venue: NeurIPS 2025 poster ()
Keywords: Operator Fusion
LLM Inference
Inter-core Interconnect
Primary area: infrastructure

Abstract: Large language model (LLM) decoding suffers from high latency due to fragmented execution across operators and heavy reliance on off-chip memory for data exchange and reduction. 
This execution model limits opportunities for fusion and incurs significant memory traffic and kernel launch overhead.
While modern architectures such as NVIDIA Hopper provide distributed shared memory and low-latency intra-cluster interconnects, they expose only low-level data movement instructions, lacking structured abstractions for collective on-chip communication.
To bridge this software-hardware gap, we introduce two cluster-level communication primitives, ClusterReduce and ClusterGather, which abstract common communication patterns and enable structured, high-speed data exchange and reduction between thread blocks within a cluster, allowing intermediate results to be on-chip without involving off-chip memory.
Building on these abstractions, we design ClusterFusion, an execution framework that schedules communication and computation jointly to expand operator fusion scope by composing decoding stages such as QKV Projection, Attention, and Output Projection into a single fused kernels.
Evaluations on H100 GPUs show that ClusterFusion outperforms state-of-the-art inference frameworks by $1.61\times$ on average in end-to-end latency across different models and configurations.

ClusterFusion: Expanding Operator Fusion Scope for LLM Inference via Cluster-Level Collective Primitive
Xinhao Luo
Zihan Liu 12-
Yangjie Zhou
Shihan Fang'
Ziyu Huang Yu Feng
Chen Zhang!
Shixuan Sun!
Zhenzhe Zheng' Jingwen Leng 1.2 Minyi Guo
Shanghai Jiao Tong University
Shanghai Qi Zhi Institute
National University of Singapore
[Ix6666, altair.liu, fang-account, huang_ziyn, wtw.yyyyyuuuueellloo.oo @sjtu.edu.cn
ly-feng, chenzhang.sjjp sunshixuan, zhengzhenzhe, @ejju.edu.cn yj yi_zhou mus.edu.sg leng-jw hlswwes.se wwt.wc.ssst.ed..co.. guo-my wew.nnye..ss..uu......
Abstract
Large language model (LLM) decoding suffers from high latency due frag- mented execution across operators and heavy reliance on off-chip memory for data exchange and reduction. This execution model limits opportunities for fusion and incurs significant memory traffic and kernel launch overhead. While modern architectures such NVIDIA Hopper provide distributed shared memory and low-latency intra-cluster interconnects, they expose only low-level data movement instructions, lacking structured abstractions for collective on-chip communication. To bridge this software-hardware gap, we introduce two cluster-level communica- tion primitives, ClusterReduce and ClusterGather. which abstract common communication patterns and enable structured, high-speed data exchange and re- duction between thread blocks within cluster, allowing intermediate results be on-chip without involving off-chip memory. Building these abstractions, we design Cluster tertiinnnoon, an execution framework that schedules communication and computation jointly to expand operator fusion scope by composing decoding stages such as QKV Projection, Attention, and Output Projection into a single fused kernels. Evaluations on H100 GPUs show that ClusterFusion outper- forms state-of-the-art inference frameworks by 1.61 x on average end-to-end latency across different models and configurations. The source code is available at stttp://ggithbb cow. xiwhaoiinohoioee luooellu///a ClusterFusinn
Introduction
Large language models (LLMs) have become a cornerstone of modern artificial intelligence sys- tems. Their applications span natural language processing 49 255. code generation [7 42), and LLM inference typically involves two stages: prefilling phase that encodes the input prompt and decoding phase that generates output tokens auto-regressively. As sequence length increase and model sizes grow, the decoding phase dominates overall inference latency, making the primary bottleneck in real-time LLM applications.
mathematical reasoning 45 54
To accelerate decoding, recent research has explored various optimizations 10
19 26
13 14 26). A growing body of research focuses optimizing execution dataflow [00 which refers to the organization of computation and communication across the parallel and memory hierarchy of modern GPUs (34). Thread blocks serve the fundamental execution units, each responsible for processing a portion of the data and typically assigned to different hardware unit such
19
55
22).
"Corresponding authors.
39th Conference on Neural Information Processing Systems (NcurIPS 2025).

---

Decoding
Latenny (ind)
Prefilling
QKV Proj
Output Proj
Attention
FFN
50
99.5%
98. 98.6%
1k
4K
16k
Sequence Length
Typical Transformer Block the fundamental component of the modern LLMs.
Figure 1:
Figure 2: Latency comparison of prefill- ing and decoding for 256 tokens.
as streaming multiprocessor (SM). However, most existing systems treat thread blocks as independent execution units. Inter-block dependencies are resolved by materializing intermediate results global memory, resulting in frequent off-chip transfers, redundant synchronization, and limited operator
fusion scope 63 53].
Recent GPU architectures offer new opportunities to address these limitations. NVIDIA Hopper introduces thread block clusters and distributed shared memory (DSMEM), which enable direct on-chip communication among blocks within the same cluster 33 Despite this potential, two major challenges remain. First, current architectures expose only low-level data movement instructions without providing high-level structured communication abstractions. Second, as our analysis in Sec. 2.3 shows, DSMEM performance highly sensitive to cluster configuration. These factors make it challenging to integrate DSMEM into real-world LLM systems.
To address these limitations, we propose two cluster-level collective primitives: Cluster and ClusterGather, which abstract common communication patterns such as reduction and aggregation. These primitives enable structured intra-cluster coordination, allowing intermediate results to be shared and combined on-chip without global memory access. Built upon these primitives, our key insight is to treat each thread block cluster as a fundamental parallel unit, using cluster-level collective communication primitives to resolve inter-block dependencies efficiently. Guided by the key insight, we propose the cluster-centric dataflow and develop ClusterFusion an execution framework that jointly schedules computation and communication to expand operator fusion scope.
Evaluation NVIDIA H100 GPUs shows that ClusterFusion achieves 1.61 speedup average in end-to-end latency compared to state-of-the frameworks. These performance gains hold across diverse model architectures (c.g.. Llama [5]], DeepSeck (II) and configurations, demonstrating the generality and effectiveness of our approach.
In summary, have made the following contributions:
We analyze the communication patterns and fusion scope in existing LLM decoding workflows, identifying fragmented kernel execution and off-chip synchronization as key barriers to efficient decoding fusion. We further profile the DSMEM mechanism on NVIDIA Hopper GPUs, revealing its potential support low-latency inter-block communication and reduce off-chip memory dependency.
We propose cluster-level collective primitives, ClusterReduce and ClusterGather, to support structured inter-block collective communication. These primitives abstract reduction and aggregation over DSMEM and enable efficient coordination between thread blocks.
We develop ClusterFusion, an execution framework that expands operator fusion via our pro- posed primitives. ClusterFusion integrates structured intra-cluster communication into the cluster- centric dataflow, fusing QKV Projection, Attention and Output Projection. This enables coordinated computation and communication without off-chip memory traffic and outperforms SOTA frameworks.
2 Background
2.1 LLM Inference Workflow and Bottlenecks
LLMs are commonly built on Transformer-based decoder-only architectures (52 As shown in Fig. each Transformer Block comprises QKV Projection, Attention, Output Projection and fordardrward network (FFN). For an input sequence X, each Attention head computes projections Q XWQ,, K; XWK. V, XWv, followed by scaled dot- dot-product Attention and Output

---

Q/K/V Projection
Attention
Output Projection
Hidden
Q/KVV
Partial Attention ourput
Attention output
TB1
TB3
TB5
TB6
TB7
Weight O/KN
Key Cache
Value Cache
Weight O
GeMV
Softmax
Rescale
Head
Head
Head
Head
On-chip
Kernel boundary
(device, (dentce.symc))
Figure 3: Existing QKV Projection, Attention, and Output Projection dataflow graph.
Projection:
Q1kj vdz
QnKy vdl
== Concat
Softmax
V1. Vi...., Softmax
Vn
Wo
(1)
Here, dg is the dimension of each head's 's query (Q) and key (K). The FFN module applies three linear layers with non -onneinear activation in between, formulated as:
FFN W3 (o(W (2) W22)
where a is non-linear activation (e.g., GELJU,, and @ denotes element-wise multiplication.
During inference, the model proceeds in two distinct stages: prefilling and decoding. In the decoding stage, the model generates tokens one time in auto-regressive manner, reusing the KV cache while appending new entries. This sequential decoding pattern inherently limits parallelism and dominates inference latency longer context lengths. In Fig. 2 the percentages represent the proportion of latency on decoding stage for different sequence lengths. Decoding accounting for over 95% of the total latency for 256 tokens sequence generation, as measured using SGLang (58), a state- of-the-art LLM serving framework. This makes decoding the dominant computational bottleneck during inference and natural target for system-level performance optimization.
2.2 Existing Communication Patterns and Fusion Scope
Numerous studies have proposed various techniques & 19 55 57] to improve the Trans- former performance. Among them, increasing attention has been given to the design of execution dataflow, which refers to the structured organization of computation stages and their associated data movement, including partitioning, scheduling, and communication, over the parallel execution model and memory hierarchy (34] decoding workloads, the dataflow plays a central role in determining end-to-end latency by governing operator scheduling and intermediate data exchange.
In existing GPU-based decoding dataflows, thread blocks commonly serve the fundamental unit of parallel execution and scheduling. Fig. iillsstrates the dataflow of Llama2-7B [51] decoding phase in SGLang [58] covering three stages: QKV Projection, Atlention, and Output Projection. Within each kernel, thread blocks are assigned to individual Attention heads and operate on disjoint tiles of the hidden dimension and KV sequence. In the OKV Projection stage, each thread block performs linear transformation on its assigned hidden states to produce local Q K, and V vectors. These outputs are written to global memory. The Attention stage is implemented with FlashDecoding each block computes partial Attention result using its corresponding Q and segmented KV cache first. separate rescaling kernel then aggregates these partial results across blocks. The final Outpur Projection is executed block-wise on the aggregated Artention output. Existing decoding dataflows exhibit two forms of inter-block communication: exchanging intermediate data (e.g., Q/K/V vectors) needed by multiple blocks, and reducing partial results across blocks (e.g., Attention output). These dependencies are resolved via off-chip memory and explicit kernel boundaries, leading to global synchronization barriers and hindering operator fusion (63 53].
19 55]. where

---

This block-isolated execution structure introduces launch overhead, off-chip memory round and global synchronization barriers between kernels. Due the absence of structured communication across thread blocks, intermediate results must be materialized to global memory and reloaded by subsequent kernels, limiting opportunities for effective on-chip data reuse and broader fusion. To overcome these limitations, we require an execution mechanism that enables collective scheduling and communication across thread blocks.
2.3
Cluster-Level Opportunities and Challenges
2-18 akooko
Compune
Cluater
Compute
Compute
Latency (cyotes)
470
Global Momory
Bandwith (TBA))
SM
130
200
120
Streed
2.5
Ski-to SM NoC
160
2.0
100
Rossttuted Shared Mamery)
16
16
16
La Good
Clustor Size
Clustor Size
Cluster Sizo
Figure NVIDIA Hop-
Figure 5: SM-to-SM access latency (left), bandwidth (middle), and number of active SMs (right) for varying cluster size on H100 GPU.
per architecture.
The previous section shows that the mainstream methods resolve inter-block data dependencies via global memory synchronization. However, recent architectures, such NVIDIA Hopper [33 shown in Fig. 4 can group a set of thread blocks into a cluster. Within each cluster, thread blocks can share data directly through high-speed SM-to-S AMM Network-on -nnhhpp (NOC), also referred to as DSMEM, thereby avoiding costly global memory accesses and enabling efficient intra-cluster communication.
To understand the opportunities and challenges of this mechanism, profile DSMEM an NVIDIA H100 GPU by varying the cluster size from to 16, which the maximum supported by Hopper hardware. As shown in Fig. [ left, SM-to-SM access latency improves significantly with small cluster sizes. When the cluster size is 2. the average latency reaches 190 cycles, which is substantially lower than global memory latency (exceeding 470 cycles). This demonstrates the potential of DSMEM for low-latency on-chip communication.
However, this benefit comes with notable trade-offs. As the cluster size increases, the available communication bandwidth decreases, slightly lagging behind the global memory bandwidth when the cluster size reaches 16 (2.90 TB/s vs. 2.96 TB/s) due crossbar architecture B]. Additionally, the number of active SMs reduced due to hardware constraints. These effects limit the scalability of cluster-based execution and necessitate careful configuration to balance communication efficiency and overall parallelism.
Beyond hardware-level trade-offs, there are software-level challenges well. NVIDIA currently exposes DSMEM and thread block cluster functionality only through low-level PTX instructions, which only support basic, peer-to-peer data movement between thread blocks K3 This low-level interface presents significant challenges for expressing reusable synchronization and communication patterns in practical scenarios, which results in a steep programming barrier and leaves developers without clear guidance on how to effectively apply these capabilities.
ClusterFusion: Execution Framework with Cluster-Centric Dataflow
This section presents the design of ClusterFusion. We begin by introducing cluster-level collective primitives that enable structured data reduction and aggregation across thread blocks. These primi- tives form the foundation for the cluster-centrie dataflow that fuses QKV Projection, Attention Output Projection. ClusterFusi builds on this to achieve high-performance LLMM decoding with expanded operator fusion scope.
Artentioon,Ateenton and
3.1 Cluster-Level Collective Primitives
To explore hardware-level trade-off's and overcome software-level challenge, we introduce two cluster- level collective primitives that abstract common communication patterns. The key insight is to treat thread block cluster as a fully connected logical network, where blocks participate in structured collective operations. Inspired by communication primitives in distributed systems [48 4 design

---

Algorithm ClusterReduce over DSMEM Thread Block view
Require: A cluster of N 2 thread blocks (A Σ 4), each with rank & [0, N 1] and shared memory buffer D2 containing local data that needs
reduced together with data from other thread blocks in the same cluster. A reduction operator @ (c.g., sum or max).
1: Allocate shared memory buffer By with the same size of D3
Used as temporary buffer.
2: stride <- 1.
3: while stride N do
4: block rank send_to <- (b stride) mod N.
5: block rank recv_frod + (b stride + N) mod N.
6:
Send D2 to Ba of block to via DSMEM.
7:
Receive Deev_trvo from block from into B, via DSMEM.
8:
Wait for the arrival of Dzecr_trox.
9: Do f- D3 + Bg
Aggregate partial result using reduction operator.
10: stride +- stride x
b Exponential stride progression.
11:
end while
12:
Return Da-
Algorithm 2 ClusterGather over DSMEM Thread Block view
Require: A cluster of N = 2* thread blocks (k < 4), cach with rank b E [0, N 1] and shared memory buffer D, of size N x size. The first segment Dz(0 size] contains the local data that needs to be gathered together with data from other thread blacks in the same cluster.
stride 1.
2: while stride N do
3: block rank send_to + stride) mod N.
4: block rank recv_from ffrommomm-- (b stride N) mod N.
5: Send Da(0: size stride] D to /strride size stride aize of block send_
via DSMEM.
6: Receive D Dreev_fron( size x stride] from block recv fromcmffrrmmoio into D, (stridd x xsize: 2 2xstrideX
size] via DSMEM.
7:
Wait for the arrival of D reex_ size stride].
8: stride <- stride x
Exponential stride progression
9:
end while
10:
Return Do-
two cluster-level primitives: ClusterReduce, which reduces data across blocks using associative operators such sum or max, and ClusterGather, which replicates local data from each block to all others for data sharing as shown in Fig.
As shown in Alg. I and Alg. 2 both ClusterReduce and ClusterGather adopt binary- pattern across log2N rounds. where N is the cluster size. In each round, the communication stride doubles, and each block exchanges data with peer whose index is offset by the current stride. ClusterReduce performs element-wise reductions while keeping the message size constant, whereas ClusterGather progressively accumulates remote data by doubling the message size in each round. This shared structure facilitates uniform implementation and hardware tuning, enabling efficient cluster-level synchronization and data sharing.
Cluster
Cluster
ArB--CHD
ArBiC.D Block
Blook
Blook
ClssteRReduce
Blook
Block
Blook
ChaterGather
Block
Block
Black A+B+CHD
Figure 6: Illustration of cluster-level collective communication primitives: ClusterReduce and ClusterGather.
5

---

Hidden
Cluster Getter
Fused Fused_QKV LKV__QVV_P__________ QKV__Koj__ATTT Proj_Attn__On Out Proj
Cluster
Reduce
T80
TB1
TB2
T83
TB4
TB5
TB6
TB7
Weight Q/K/V
Key Cache
Value Cache
Weight O
Softmax
Rescale
Head
Head
On-chip
Kernel boundary sy..smm))
Figure
Cluster-centric fused QKV Projection, Attention and Output Projection dataflow graph.
3.2 Cluster-Centric Dataflow Design with Primitives
Building upon the proposed cluster-level primitives, we now illustrate how they are employed to construct cluster-centric dataflow with expanded operator fusion scope. The core idea is to treat the thread block cluster as cooperative execution and scheduling unit. Data-dependent dimensions kept within each cluster resolve inter-block data dependencies using cluster-level collective primitives which avoid off-chip data exchange, while iata-independent dimensions are distributed across clusters. This organization aligns computation with memory locality and enables seamless kernel fusion by allowing intermediate results to be reused entirely within on-chip memory.
In the decoding stage. the dataflow for the Projection and Attention modules is parallelized over multiple dimensions. Specifically, the Projection is parallel along the number of heads and head dimension, while the Attention is parallel over the number of heads and the KV cache sequence length. Among these dimensions, thread blocks that compute different partitions of the head dimension and the KV cache sequence length exhibit inter-block data dependencies, since each block computes either a partial result that requires reduction or a segment result that needs to be gathered to form the final output. These thread blocks be grouped into cluster resolve data dependencies on-chip by using cluster-level collective primitives. Since Attention heads are independent across all three modules, each cluster is accordingly mapped to single head. As shown in Fig. 7 under this design, the intermediate results produced by the QKV Projection naturally remain on chip and are directly reused by the Artention module. Likewise, the output of the Attention module stays on chip and is immediately consumed by the Output Projection, enabling seamless data reuse across three modules.
By leveraging cluster-level communication primitives, implement the fused QKV Projection, Attention and Output Projection dataflow. As illustrated in Alg. 5 the dataflow is parallel in the number of heads. Each head corresponds block is assigned a rank b E N 1]. Within each cluster, thread blocks respectively partition the head dimension in QKV Projection, the KV cache token dimension in Attention, and the output dimension in Output Projection. For the whole dataflow, each thread block processes the entire input hidden states and computes the corresponding output tile O₂ after the Output Projection. In this algorithm, B denotes the batch size, D the input hidden dimension, H the total head dimension, and h, s and d represent the partitioned sizes of the head dimension, sequence length and the output dimension, thread block, respectively.
cluster of N = 2 thread blocks (k: 4), where each
cluster-centric dataflow design principle, also propose several dataflow variants, which are presented the Appendix B. To evaluate these variants, we conduct a quantitative analysis of DSMEM traffic [7] 59). We begin by analyzing the DSMEM memory traffic incurred by the ClusterReduce and ClusterGather primitives as follows:
According
Traf ficneduce (size, N) = sizexlog, NxN, ficgtherr N) = sizex (2 N +1 -1)
1) XN
Here, Traʃ Praffficeede ficReduce and Traf prafficcaata fiCGather denote the DSMEM traffic for ClusterReduce and ClusterGather, respectively. The variable size represents the size the shared memory buffer D3 in Alg. as well as the initial segment size in Alg. 2 Based on this analytical model, we estimate

---

Algorithm 3 Fused QKV Projection, Atlention and Oulput Projection Dataflow Thread Block View
Require: Input hidden states H E RIXD, QKV Projection weight W E Px3h, Output Projection weight W E R fxx and KV cache Kymcee Vi E Rax H in global memory.
b
1: Allocate shared memory buffers: Qo, Ko. Vo E Rixh, and Sum: Smax (softmax statistics).
2: Compute segment results of QKV Projection: Qs, K, V H, X WQKV
3: Obatin the complete QKV: (Qr, Ko, Vo) ClusterGather Ko: Vo).
4: Compute partial result of Attention similar to the FlashDecoding dataflow:
Compute S, <- exp( Q6 (Kguche Ks)²), obtain local Sam, Smax.
And store Smax in register Reg max Compute A, S₁ (Vguche, V&b
softmax statistics.
>> As reuse the shared memory space of Qs
5: Obtain the complete softmax statistics Sam and Smax:
Sum f- ClusterReduce( Sum sum), Smax C--ssstereeduce (Smxx max)
6: Rescale Attention output according to online softmax:
A Ab exp( Reg S
Saum
7: Obatin the complete Attention output: AB +- ClusterReduce(A As, sum).
8: Compute the result the Output Projection and write it to global memory using atomi cAdd:
Og + As x A
9: return Og-
the total DSMEM memory traffic of the dataflow used in Alg. which includes one ClusterGather and two ClusterReduce operations:
Traf ficTora f ficReduce (3h, + Traf (H,
We omit the traffic generated by the softmax statistics, it involves only floats and is negligible compared to tensor. According to the analysis shown Appendix B, this dataflow yields the lowest DSMEM traffic and achieves better performance compared other dataflow variants.
The proposed cluster-centric dataflow design principle can be naturally generalized to other operators, including DeepSeek MLA [II] The corresponding algorithms provided in the Appendix B. We implement an end-to-end execution framework, ClusterFusion, based on the cluster-centric dataflow, which incorporates the aforementioned fused QKV Projection, Attention, and Output Projection modules. For other components such as FFN and RMSNorm, we adopt optimized imple- mentations consistent with those in existing frameworks such CUTLASS 36) and Flashinfer (55)
4 Evaluation
Experimental Setup We evaluate ClusterFusion on an NVIDIA H100 SXM5 80GB GPU [3 For the end-to-end evaluation, all inputs are in FP16 precision, and the context length varies from IK to 16K. We set the batch size to I; results of multi batch are presented the Appendix C. All experiments are conducted using PyTorch 2.5.1 40 and CUDA 12.4
34
Baselines We compare ClusterFusion with four state-of-the-art LLM inference frameworks: SGLang 0.4.3.post2 (58), vL..M 0.6.4.postl
(24), TensorRTT-LLM 0.18.0 39), and MLC-LLM
0.20.dev0 [32). For all baselines, we use the recommended configurations from their official documen- tation, including backend kemnlss from librarics such CUTLASS (36), FlashAttention 9 8 44 57].
and FlashInfer (55), or generated by Triton [50] and TVM [7]. All frameworks enable CUDA
Graph 335 and Torph.compile [5]
Models ClusterFusion is evaluated on two representative LLMs: Llama2-7B (51 and DeepSeek- V2-Lite [1]. both based on the Transformer architecture. Llama2-7B adopts standard Multi-Head Attention (MHA) mechanism, while DeepSeek- V2-Lite employs Multi-head Latent Attention (MLA) algorithm. These models differ in hidden dimensions, head dimensions, and number of heads.
7

---

SGLang YLLM TRT-LLM
SGLang
YLLM
TRE-LLM
ClusierFusion
CluslerFusion
300
TPOT (ms)
Latency (us)
250
10
150
100
50
16k
16k
Sequence Length
Sequence Length
Sequercee Length
Sequence Length
Figure 8: TPOT of Llama2-7B (left) and DeepSeck- SeepSeck-V2-Lite (right) on H100.
Figure 9: Latency of core modules of Llama2-78 -BB (left) and DeepSeek- -Vt-Lite (right) on H100.
4.1 End-to-End and Core Module Evaluation
We use time per output token (TPOT) as the metric for end-to-end evaluation. The results are presented in Fig. 17 For baselines, the results include both CUDA graph launch overhead and kernel execution latency. On average, ClusterFusion achieves 1.41x, 1.39x, 1.43x, x, and x.03x speedups over SGLang, vLLM, TensorRT-LLM, and MLC-LLM across various sequence lengths with a cluster size of on Llama2-7B. For DeepSeck- V2-Lite, ClusterFusion delivers average speedups of x.34x, 1.37 x, 1.51 x, and 2.39x under the same conditions.
As shown in Fig. 18 for the core QKV Projection, Attention, and Ouiput Projection modules, ClusterFusion achieves average speedups of 1,85x, 1.73x, 1.61 x, and 3.19x compared to SGLang, VLLM, TensorRT-LLLM, and MLC-LLM, respectively on Llama2-7B. Similarly, for DeepSeek-V2-1 V2-Liite, ClusterFusion delivers average speedups of 1.66x, 1,64x, x.35x, and 3.5 For the DeepSeek MLA, which is specifically designed to better leverage GPU hard- ware, the optimization space for operator fusion relatively lim- ited. Nevertheless, as shown in Fig. 10 sequence lengths in real datasets are predominantly under 8k. In this range, ClusterFusion achieves significant performance improvements, demonstrating its effectiveness in real- world scenarios.
Percentage (%)
1.. ShareGPT Splitwie
40
20
0
5000
0 5000
Sequence Length
Figure 10: distribution in ShareGPT [
Sequence length
and Splitwise [41 2 datasets.
We further evaluate the performance of the core modules in ClusterFusi on under varying cluster sizes and numbers of Altention heads, with sequence lengths of 4K and 16K. The results are presented in Fig. II In our design, each Attention head is mapped to cluster, so the cluster size determines the internal parallelism within each head. When the number of Attention heads is 32 and 64, cluster size of 4 yields the best performance. However, when the number of heads increases to 128, smaller cluster size of 2 becomes optimal. Conversely, cluster sizes of and 16 lead worse performance due to increased interconnect latency, bandwidth contention, and reduced number of active SMs, which collectively limit overall core utilization, as illustrated in Fig. S Based on both theoretical analysis and empirical evidence, we conclude that the optimal cluster size varies across workloads. Therefore, cluster size should be tuned accordingly to maximize performance.
4.2 Speedup Analysis
We identify two primary factors contributing to the performance ad advantage of ClusterFusion over existing frameworks with CUDA Graph optimizations: minimized global memory transfer size and
Seq length 4k
Seq length
18k
Wenoryy Transtee Sioe (M00)
SGLing SLLM
16
au3
Aansa3-71
DeepSeek-V--Liee
Lourch Ovelicud
ChoterFiddrr
Cluster
8
124
285
500
1.0-
:00
013
221
177
197
390
5003
0.5
143
320
32
84
128
32
64
128
400
1k
24
200
Llanna
Doop
Head Num
Seiguenor Length
Models
Figure 11: Latency of core module in ClusterFusion with varying settings.
Figure 12: Comparison global memory data transfer size (left) and GPU kernel launch overhead (right).
8

---

To quantify these benefits, we leverage NVIDIA Nsight Systems 38 and Nsight Compute [37 to profile global memory transfer volume and GPU kernel launch overhead across different models and configurations. The results are illustrated in Fig. 19 The performance gain stems from the fact that ClusterFusion executes QKV Projection, Attention, and Ouiput Projection entirely on-chip, significantly reducing intermediate memory traffic. Additionally, ClusterFusion reduces kernel launch overhead by nearly an order of magnitude in end-to-end scenarios, even when compared to baselines optimized with CUDA Graph.
reduced GPU kerne launch overhead 63 53
4.3 Additional Analysis
Table I: Latency comparison of on-chip ClusterReduce and ClusterGather with DSMEM versus off-chip imple- mentations without DSMEM
12-
WHO DSMEM
Wi DSMEM
TPO7 (mss)
10
Operation
Data Size (KB)
Off-chip (1x))
Ou-chip (ae)
Speedup
8.03
6,77
1.18x
1.36x
ClusterRedace
128
14.95
7.42
2.01x
22.44
9.17
2.44x
16k
6,26
3.90
1.00x
Sequence Length
ClusterGather
6.27
4.39
4,12
1.52x
1.44x
Figure 13: TPOT of ClusterFusion
256
6.61
4.15
1.59
Llama2-7B with and without DSMEM.
To demonstrate the importance of the on-chip interconnect leveraged in our dataflow design, we conduct microbenchmark to evaluate the cluster-level collective communication primitives intro- duced in Sec. 3.1 As shown in Tbl. [ the on-chip ClusterReduce and ClusterGather operations over DSMEM exhibit significantly lower latency across varying data transfer sizes compared the off-chip implementations.
We further perform ablation studies compare ClusterFusion with and without DSMEM enabled. The experimental results are presented Fig. [3] Across different sequence lengths, disabling DSMEM increases the time per output token (TPOT) by up 33%. These results highlight the effectiveness of DSMEM and the cluster-level collective primitives in enabling efficient on-chip reduction and aggregation, thereby improving end-to-end inference performance.
Discussion on Fusion Scope and Architectural Outlook
ClusterFusion builds on intra-cluster DSMEM communication, where each fused scope is bounded by fixed cluster size (up to 16 thread blocks) 33). This imposes constraints fusion granularity and scheduling ficxibility. Although most decoding operators in today's 's mainstream LLMs [T such Projection and Attention, fit comfortably within this limit, future models with larger hidden dimensions or specialized operator variants may challenge this boundary. When fused operators exceed the cluster scope, the system must fall back to global memory communication. introducing additional latency and runtime fragmentation. This motivates reflection on hardware support for broader intra-chip collectives 20 27 18).
Our findings suggest that enabling low-latency, topology- communication across a broader set of SMs would unlock more uniform and scalable fusion strategies. Such architectural support could extend structured coordination beyond current cluster boundaries. ClusterFusion highlights this co-design opportunity practical step toward supporting the growing complexity of architectures.
Conclusion
This paper presents ClusterFusion, an execution framework that schedules communication and computation jointly to expand operator fusion scope by composing operators during decoding stages such as QKV Projection, Attention, and Output Projection into single fused kernels. By incorporating cluster-level collective communication primitives, ClusterFusic on effectively reduces global memory traffic and kernel launch overhead, enabling efficient on-chip execution of key L.M decoding modules. Our comprehensive evaluation NVIDIA H100 GPUs with representative

---

models Alama-2-7B 7B and DeepSeck-V2-Lite V2-Lite demonstrates that ClusterFusion outperforms state-of- the-art LLM inference frameworks across different models and configurations.
Acknowledgement
This work was supported by the National Key R&D Program of China under Grant 2022YFB4501400, and the National Natural Science Foundation of China (NSFC) grant 62222210. This work was also supported by Shanghai Qi Zhi Institute Innovation Program SQZ202316. We sincerely thank Yilu Huang and Chiheng Jin for their assistance with kernel implementation and end-to-end integration. We would also like to thank the anonymous reviewers, as well as the Area Chairs and Program Chairs, for their constructive feedback and valuable comments that helped improve this work.
References
[1] Sharegpt.https:// httpeppi.httpps/ssgggptt..o
[2] Splitwise.
[3]
[4] Reza Yazdani Aminabadi, Samyam Rajbhandari, Ammar Ahmad Awan, Cheng Li, Du Li, Elton Zheng, Olatunji Ruwase, Shaden Smith, Minjia Zhang. Jeff Rasley, and Yuxiong He. Deepspeed-inferee inference: Enabling efficient inference of transformer models at unprecedented scale. In Felix Wolf, Sameer Shende, Candace Culhane, Sadaf R. Alam, and Heike Jagode, editors, SC22: International Conference for High Performance Computing, Networking, Storage and Analysis, Dallas, TX, USA, November 13-18, 2022, pages 16:1-46:5. IEEE, 2022.
[5] Jason Ansel, Edward Z. Yang, Horace He, Natalia Gimelshein, Animesh Jain, Michael Voznesen- sky, Bin Bao, Peter Bell, David Berard, Evgeni Burovski, Geeta Chauhan, Anjali Chourdia, Will Constable, Alban Desmaison, Zachary Vito,, Elias Ellison, Will Feng, Jiong Gong, Michael Gschwind, Brian Hirsh, Sherlock Huang. Kshiteej Kalambarkar, Lauuent Kirsch, Michael Lazos, Mario Lezcano. Yanbo Liang, Jason Liang, Yinghai Lu, C. K. Luk. Bert Maher, Yunjie Pan, Christian Puhrsch, Matthias Reso, Mark Saroufim, Marcos Yukio Siraichi, Helen Suk, Shunting Zhang, Michael Suo, Phil Tillet, Xu Zhao, Eikan Wang, Keren Zhou, Richard Zou, Xiaodong Wang. Ajit Mathews, William Wen, Gregory Chanan, Peng Wu, and Soumith Chintala. Pytorch 2: Faster machine learning through dynamic python bytecode transformation and graph pilation. In Proceedings of the 29th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 2, ASPLOS 2024, La Jolla, CA. May 2024, pages 929-947. ACM, 2024.
USA, 27 April 2024-
[6] Cerebras. Cerebras wse-3 chip. https: http:::/wwwwwweeebbbea ailesaiiiiciii ai/chip
[7] Tianqi Chen, Thierry Moreau, Ziheng Jiang, Lianmin Zheng, Eddie Q. Yan, Haichen Shen, Meghan Cowan, Leyuan Wang, Yuwei Hu, Luis Ceze, Carlos Guestrin, and Arvind Krishna- murthy. TVM: an automated end-to-end optimizing compiler for deep learning. 13th USENIX Symposium Operating Systems Design and Implementation, OSDI 2018, Carlsbad, CA. USA, October 8-10, 2018, pages 578-594. USENIX Association, 2018
[8] Tri Dao. Flashattention-2: Faster attention with better parallelism and work partitioning. In The Twelfth International Conference on Learning Representations, ICLR 2024, Vienna, Austria, May 7-11, 2024. OpenReview.net, 2024.
[9] Tri Dao, Daniel Y. Fu, Stefano Ermon, Atri Rudra, and Christopher Ré. Flashattention: Fast and memory-effiiient exact attention with io-awareness. In Advances in Neural Information Processing Systems 35: Annual Conference on Neural Information Processing Systems 2022,
NeurlPS 2022, New Orleans, LA, USA, November 28 December 9, 2022, 2022.
[10] Tri Dao, Daniel Haziza, Francisco Massa, and Grigory Sizov. Flash-decoding for long-context inference. httph::////// //cts//////....// ora
html 2023.
10

---

[11] DeepSeek-AI, Aixin Liu, Bei Feng, Bin Wang, Bingxuan Wang, Bo Liu. Chenggang Zhao, Chengqi Deng, Chong Ruan, Damai Dai, Daya Guo, Dejian Yang, Deli Chen, Dongjie Ji, Erhang Li, Fangyun Lin, Fuli Luo, Guangho Hao, Guanting Chen, Guowei Li, Hao Zhang, Hanwei Xu, Hao Yang. Haowei Zhang. Honghui Ding. Huajian Xin, Huazuo Gao, Hui Li, Hui Qu, J. I. Cai, Jian Liang, Jianzhong Guo, Jiaqi Ni, Jiashi Li, Jin Chen, Jingyang Yuan, Junjie Qiu, Junxiao Song, Kai Dong, Kaige Gao, Kang Guan, Lean Wang. Lecong Zhang, Lei Xu, Leyi Xia, Liang Zhao, Liyue Zhang, Meng Li, Miaojun Wang. Mingchuan Zhang, Minghua Zhang, Minghui Tang. Mingming Li, Ning Tian, Pampan Huang. Peiyi Wang. Peng Zhang, Qihao Zhu, Qinyu Chen, Qiushi Du, R. J. Chen, R. L. Jin, Ruiqi Ge, Ruizhe Pan, Runxin Xu, Ruyi Chen, S. S. Li, Shanghao Lu, Shangyan Zhou, Shanhuang Chen. Shaoqing Wu, Shengfeng Ye, Shirong Ma, Shiyu Wang. Shuang Zhou, Shuiping Yu, Shunfeng Zhou, Size Zheng, Tao Wang, Tian Pei, Tian Yuan, Tianyu Sun, W. L.. Xiao, Wangding Zeng. Wei An, Wen Liu, Wenfeng I Liang, Wenjun Gao, Wentao Zhang, X. Q. Li, Xiangyue Jin, Xianzu Wang, Xiao Bi, Xiaodong I Liu, Xiaohan Wang, Xiaojin Shen, Xiaokang Chen, Xiaosha Chen, Xiaotao Nie, and Xiaowen Sun. Deepseek-v2: A strong, economical, and efficient mixture-of-experts language model. CoRR, abs/2405.044333 2024.
[12] DeepSeek-AI, Aixin Liu, Bei Feng, Bing Xue, Bingxuan Wang, Bochao Wu, Chengda Lu, Chenggang Zhao, Chengqi Deng, Chenyu Zhang. Chong Ruan, Damai Dai, Daya Guo, Dejian Yang, Deli Chen, Dongjie Ji, Erhang Li, Fangyun Lin, Fucong Dai, Fuli Luo, Guangho Hao, Guanting Chen, Guowei Li, H. Zhang, Han Bao, Hanwei Xu, Haocheng Wang. Haowei Zhang, Honghui Ding, Huajian Xin, Huazuo Gao, Hui Li, Hui Qu, J. L. Cai, Jian Liang, Jianzhong Guo, Jiaqi Ni, Jiashi Li, Jiawei Wang, Jin Chen, Jingchang Chen, Jingyang Yuan, Junjie Qiu, Junlong Li, Junxiao Song, Kai Dong. Kai Hu, Kaige Gao, Kang Guan, Kexin Huang, Kuai Yu, Lean Wang, Lecong Zhang. Lei Xu, Leyi Xia, Liang Zhao, Litong Wang, Liyue Zhang. Meng Li, Miaojun Wang, Mingchuan Zhang, Minghua Zhang. Minghui Tang, Mingming Li, Ning Tian, Panpan Huang, Pciyi Wang. Peng Zhang, Qiancheng Wang. Qihao Zhu, Qinyu Chen, Qiushi Du, J. Chen, R. I. Jin, Ruiqi Ge. Ruisong Zhang, Ruizhe Pan, Runji Wang, Runxin Xu, Ruoyu Zhang, Ruyi Chen, S. S. Li, Shanghao Lu, Shangyan Zhou, Shanhuang Chen, Shaoqing Wu, Shengfeng Ye, Shengfeng Ye, Shirong Ma, Shiyu Wang. Shuang Zhou, Shuiping Yu, Shunfeng Zhou, Shuting Pan, T. Wang, Tao Yun, Tian Pei, Tianyu Sun, W. L. Xiao, and Wangding Zeng. Deepseek-v3 technical report. CORR, abs/2412.194377 2024.
[13] Yue Guan, Changming Yu, Yangjie Zhou, Jingwen Leng, Chao Li, and Minyi Guo. Fractal: Joint multi-level sparse pattern tuning of accuracy and performance for DNN pruning. In Rajiv Gupta, Nael B. Abu-Ghazalehh Madan Musuvathi, and Dan Tsafrir, editors, Proceedings of the 29th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 3, ASPLOS 2024, La Jolla, CA, USA, 27 April 2024- I May 2024, pages 416-430. ACM, 2024.
[14] Cong Guo, Bo Yang Hsueh, Jingwen Leng, Yuxian Qiu, Yue Guan, Zehuan Wang. Xi- aoying Jia, Xipeng Li, Minyi Guo, and Yuhao Zhu. Accelerating sparse DNN models without hardware-support via tile-wise sparsity. In Christine Cuicchi, Irene Qualters, and William T. Kramer, editors, Proceedings of the International Conference for High Performance Computing, Networking, Storage and Analysis. SC 2020, Virtual Event Atlanta, Georgia. USA. November 9-19, 2020, page 16. IEEE/ACM, 2020.
[15) Cong Guo, Yuxian Qiu, Jingwen Leng, Xiaotian Gao, Chen Zhang, Yunxin Liu, Fan Yang, Yuhao Zhu, and Minyi Guo. Squant: On-the-fly data-free quantization via diagonal hessian approximation. In The Tenth International Conference on Learning Representations, ICLR 2022, Virtual Event, April 25-29, 2022. OpenReview.net. 2022.
[16] Cong Guo, Chen Zhang, Jingwen Leng, Zihan Liu, Fan Yang, Yunxin Liu, Minyi Guo, and Yuhao Zhu. ANT: exploiting adaptive numerical data type for low-bit deep neural network quantization. In 55th IEEF/ACM International Symposium on Microarchitecture, MICRO 2022, Chicago. IL. USA, October 1-5, 2022. pages 1414-1433, IEEE, 2022.
[17] Daya Guo, Qihao Zhu, Dejian Yang. Zhenda Xie, Kai Dong, Wentao Zhang, Guanting Chen, Xiao Bi, Y. Wu, Y. K. Li, Fuli Luo, Yingfei Xiong, and Wenfeng Liang. Deepseek-coder: When the large language model meets programming the rise of code intelligence. CoRR, abs/2401. 14196, 2024.
11

---

[18] Congjie He, Yeqi Huang, Pei Mu, Ziming Miao, Jilong Xue, Lingxiao Ma, Fan Yang, and Luo Mai. Waferllm: A wafer-scale LLM inference system. CoRR, abs/2502.04563, 2025.
[19] Ke Hong. Guohao Dai, Jiaming Xu, Qiuli Mao, Xiuhong Li, Jun Liu, Kangdi Chen, Yuhan Dong, and Yu Wang. Flashdecodingtt: Faster large language model inference with asynchronization, flat GEMM optimization, and heuristics. In Proceedings of the Seventh Annual Conference on Machine Learning and Systems, MLSys 2024, Santa Clara, CA, USA, May 13-16, 2024. misys.org, 2024.
[20] Zhe Jia, Blake Tillman, Marco Maggioni, and Daniele Paolo Scarpazza. Dissecting the graphcore IPU architecture via microbenchmarking CoRR, abs/1912.0344 0231112...11... 2019.
[21] Huiqiang Jiang, Yucheng Li, Chengruidong Zhang. Qianhui Wu, Xufang Luo. Surin Ahn, Zhenhua Han, Amir H. Abdi, Dongsheng Li, Chin-Yew Lin, Yuging Yang, and Lili Qiu. Minference 1.0: Accelerating pre-filling for long-context Ilms via dynamic sparse attention, 2024.
[22] Shengyu Liu Jiashi Li. Flashmla: Efficient mla decoding kernels. https:// com/ deepseek- Siapseek-aiFFleee ai/FlashMLA 2025.
[23] Zhixian Jin, Christopher Rocca, Jiho Kim, Hans Kasan, Minsoo Rhu, Ali Bakhoda, Tor M. Aamodt, and John Kim. Uncovering real gpu noc characteristics: Implications on intercon- nect architecture. In 2024 57th IEEE/ACM International Symposium on Microarchitecture (MICRO), pages 885-898, 2024.
[24] Woosuk Kwon, Zhuohan Li, Siyuan Zhuang, Ying Sheng, Lianmin Zheng, Cody Hao Yu, Joseph Gonzalez, Hao Zhang, and Ion Stoica. Efficient memory management for large language model serving with pagedattention. In Proceedings of the 29th Symposium on Operating Systems Principles, SOSP 2023, Koblenz, Germany, October 23-26, 2023, pages 611-626. ACM, 2023.
[25] Christoph Leiter, Ran Zhang, Yanran Chen, Jonas Belouadi, Daniil Larionov, Vivian Fresen, and Steffen Eger. Chatgpt: A meta- analyss after 2.5 months. CoRR, abs/2302.13795, 2023.
[26] Lin, Jiaming Tang, Haotian Tang, Shang Yang. Wei-Ming Chen, Wei-Chen Wang, Guangxuan Xiao, Xingyu Dang. Chuang Gan, and Song Han. AWQ: activation-aware weight quantization for on-device LLM compression and acceleration. In Proceedings of the Seventh Annual Conference on Machine Learning and Systems, MLSys 2024, Santa Clara, CA, USA, May 13-16, 2024. mlsys.org, 2024.
[27] Yiqi Liu, Yuqi Xue, Yu Cheng, Lingxiao Ma, Ziming Miao, Jilong Xue, and Jian Huang, Scaling deep learning computation over the inter-core connected intelligence processor with T10. In Proceedings of the ACM SIGOPS 30th Symposium on Operating Systems Principles, SOSP 2024, Austin, TX, USA, November 4-6, 2024, pages 505-521. ACM, 2024.
[28] Zihan Liu, Jingwen Leng, Zhihui Zhang, Quan Chen, Chao Li, and Minyi Guo. VELTAIR: towards high-peeformance multi-tenant deep learning services via adaptive compilation and scheduling. In Babak Falsafi, Michael Ferdman, Shan Lu, and Thomas F. Wenisch, editors, ASPLOS $22: 27th ACM International Conference on Architectural Support for Programming Languages and Operating Systems. Lausanne, Switzerland, 28 February 2022 4 March 2022, pages 388-401. ACM, 2022.
[29] Zihan Liu, Xinhao Luo, Junxian Guo, Wentao Ni, Yangjie Zhou, Yue Guan, Cong Guo, Weihao Cui, Yu Feng, Minyi Guo, Yuhao Zhu, Minjia Zhang, Chen Jin, and Jingwen Leng. VQ- LLM: nigh--peffrmmncc code generation for vector quantization augmented LLM inference. In IEEE International Symposium on High Performance Computer Architecture, HPCA 2025, IEEE, 2025.
Las VVg USA, March 1-5, 2025, pages
[30] Zihan Liu, Wentao Ni, Jingwen Leng, Yu Feng, Cong Guo, Quan Chen, Chao Li, Minyi Guo, and Yuhao Zhu. JUNO: optimizing high nigh-dimensional approximate nearest neighbour search with sparsity-aware awarre algorithm and ray- -raciinn core mapping. In Rajiv Gupta. Nael B. Abu-Ghazaleh, Madan Musuvathi, and Dan Tsafrir, editors, Proceedings of the 29th ACM International Conference on Architectural Support for Programming I Languages and Operating Systems, Volume 2, ASPLOS 2024, La Jolla, CA, USA, 27 April 2024- 1 May 2024, pages 549- ACM, 2024.
12

---

[31] Enzhe Lu, Zhejun Jiang, Jingyuan Liu, Yulun Du, Tao Jiang, Chao Hong, Shaowei Liu, Weiran He, Enming Yuan, Yuzhi Wang, Zhiqi Huang, Huan Yuan, Suting Xu. Xinran Xu, Guokun Lai, Yanru Chen, Huabin Zheng, Junjie Yan, Jianlin Su, Yuxin Wu, Neo Y. Zhang, Zhilin Yang, Xinyu Zhou, Mingxing Zhang. and Jiezhong Qiu. Moba: Mixture of block attention for long-context llms. CoRR, abs/2502. 13109, 2025.
[32]
MLC team. MLC-LLM, 2023-2025.
[33]
NVIDA. Nvidia hopper architecture.
technologies/hoppee http://peeeeoopprrrrrrtttteeeeee architectchht
[34] NVIDIA. Cuda c++ programming guide.
https://aocs.nvicca.. /
cuda- proogramrggg wuwde//nddx.ehtmll
[35] NVIDIA. / Chtelp/....taa///eeeooomm//////r
[36] NVIDIA. Cutlass. httpss//gthhhb.cmm////DDD//aaas
[37] NVIDIA. Nvidia nsight compute. http::tttt:::dellpppeeeeo http://wwwppppar.iiiiaaa.ooaaaaa
[38] NVIDIA. Nvidia http:h/twtpp..////ooopepe/r/a//a
[39]
NVIDIA. Tensorrt-Ilm. http:://///tbbbb////////////sttt
[40] Adam Paszke, Sam Gross, Francisco Massa, Adam Lerer, James Bradbury, Gregory Chanan, Trevor Killeen, Zeming Lin, Natalia Gimelshein, Luca Antiga, Alban Desmaison, Andreas Köpf. Edward Z. Yang, Zachary Vito, Martin Raison, Alykhan Tejani, Sasank Chilamkurthy, Benoit Steiner, Lu Fang, Junjie Bai, and Soumith Chintala. Pytorch: An imperative style, high- performance deep learning library. In Advances in Neural Information Processing Systems Annual Conference on Neural Information Processing Systems 2019, NeurIPS 2019,
32:
December 8-14, 2019, Vancouver, BC, Canada, pages 8024 8024-8035, 2019.
[41] Pratyush Patel, Esha Choukse, Chaojie Zhang, Aashaka Shah, fñigo Goiri, Saeed Maleki, and Ricardo Bianchini. Splitwise: Efficient generative LLM inference using phase splitting. In st ACM/IEEE Annual International Symposium on Computer Architecture, ISCA 2024, Buenos Aires, Argentina, June 29 July 3, 2024, pages 118-132. IEEE, 2024.
[42] Baptiste Rozière, Jonas Gehring. Fabian Gloeckle, Sten Sootla, Itai Gat, Xiaoqing Ellen Tan, Yossi Adi, Jingyu Liu, Tal Remez, Jérémy Rapin, Artyom Kozhevnikov, Ivan Evtimov, Joanna Bitton, Manish Bhatr, Cristian Canton-Ferrer, Aaron Grattafiori, Wenhan Xiong, Alexandre Défossez, Jade Copet, Faisal Azhar, Hugo Touvron, Louis Martin, Nicolas Usunier, Thomas Scialom, and Gabriel Synnaeve. Code llama: Open foundation models for code. CoRR, abs/2308. 12950, 2023.
[43] SGLangg Sglang deepseek model optimizations.
sgl- 1llannnnnn gggl-leamingg-aterrial/s///// http//////mainn/iiiid/////aaaaaa deegglanneedeppsee__ooddee optimizations. pdf
[44] Jay Shah, Ganesh Bikshandi, Ying Zhang, Vijay Thakkar, Pradeep Ramani, and Tri Dao. Flashattention-3: Fast and accurate attention with asynchrony and low-precision. In Advances in Neural Information Processing Systems 38: Annual Conference on Neural Information Processing Systems 2024, NeurlPS 2024, Vancouver, BC, Canada, December 10 15, 2024, 2024.
[45] Zhihong Shao, Peiyi Wang, Qihao Zhu, Runxin Xu, Junxiao Song, Mingchuan Zhang, Y. K. Li, Y. Wu, and Daya Guo. Deepseekmath: Pushing the limits of mathematical reasoning in open language models. CoRR, abs/2402.03300. 2024.
[46] Noam Shazeer. Fast transformer decoding: One write-head is all you need. CoRR.
abs/1911. 1.02150, 2019.
[47] Yining Shi, Zhi Yang, Jilong Xue, Lingxiao Ma, Yuqing Xia, Ziming Miao, Yuxiao Guo, Fan Yang, and Lidong Zhou. Welder: Scheduling deep learning memory access via tile-graph. In 17th USENIX Symposium on Operating Systems Design and Implementation, OSDI 2023, Boston, MA, USA, July 10-12 2023, pages 701-718. USENIX Association, 2023,
13

---

[48] Mohammad Shoeybi, Mostofa Patwary, Raul Puri, Patrick LeGresley, Jared Casper, and Bryan Catanzaro. Megatron-Ime Im: Training multi-biliion parameter language models using model parallelism. CoRR, hts/1909.080553 2019.
[49] Yixuan Su, Tian Lan, Yan Wang, Dani Yogatama, Lingpeng Kong, and Nigel Collier. A contrastive framework for neural text generation. In Advances in Neural Information Processing Systems 35: Annual Conference on Neural Information Processing Systems 2022, NeurIPS 2022, New Orleans, LA, USA, November 28- December 9, 2022, 2022.
[50] Philippe Tillet, Hsiang-TTuun Kung, and David D. Cox. Triton: an intermediate language and compiler for tiled neural network computations. In Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages, MAPL@PLD. 2019, Phoenix, AZ, USA, June 22, 2019, pages 10-19. ACM, 2019.
[51] Hugo Touvron, Louis Martin, Kevin Stone, Peter Albert, Amjad Almahairi, Yasmine Babaei, Nikolay Bashlykov, Soumya Batra, Prajjwal Bhargava, Shruti Bhosale. Dan Bikel, Lukas Blecher, Cristian Canton-Ferrer, Moya Chen, Guillem Cucurull, David Esiobu, Jude Fernandes, Jeremy Fu, Wenyin Fu, Brian Fuller, Cynthia Gao, Vedanuj Goswami, Naman Goyal, Anthony Hartshomn, Saghar Hosseini, Rui Hou, Hakan Inan, Marcin Kardas, Viktor Kerkez, Madian Khabsa, Isabel Kloumann, Artem Korenev, Punit Singh Koura, Marie-Anne Lachaux, Thibaut Lavril, Jenya Lee, Diana Liskovich, Yinghai Lu, Yuning Mao, Xavier Martinet, Todor Mihaylov, Pushkar Mishra, Igor Molybog. Yixin Nie, Andrew Poulton, Jeremy Reizenstein, Rashi Rungta, Kalyan Saladi, Alan Schelten, Ruan Silva, Eric Michael Smith, Ranjan Subramanian, Xiao- qing Ellen Tan, Binh Tang. Ross Taylor, Adina Williams, Jian Xiang Kuan, Puxin Xu, Zheng Yan, Iliyan Zarov, Yuchen Zhang, Angela Fan, Melanie Kambadur, Sharan Narang, Aurélien Rodriguez, Robert Stojnic, Sergey Edunov, and Thomas Scialom. Llama 2: Open foundation and fine-tuned chat models. CoRR, abs/2307.0922 2023.
[52] Ashish Vaswani, Noam Shazeer, Niki Parmar, Jakob Uszkoreit, Llion Jones, Aidan N. Gomez, Lukasz Kaiser, and Illia Polosukhin. Attention is all you need. In Advances in Neural Information Processing Systems 30: Annual Conference on Neural Information Processing Systems 2017, December 4-9, 2017, Long Beach, CA, USA, pages 5998-6008 2017.
[53] Chunwei Xia, Jiacheng Zhao, Qianqi Sun, Zheng Wang, Yuan Wen, Teng Yu, Xiaobing Feng, and Huimin Cui. Optimizing deep learning inference via global analysis and tensor expres- sions. In Proceedings of the 29th ACM International Conference
Architectural Support for Programming Languages and Operating Systems, Volume 1, ASPLOS 2024, La Jolla, CA. USA, 27 April 2024- 1 May 2024, pages 286-301. ACM, 2024.
[54] An Yang, Beichen Zhang, Binyuan Hui, Bofei Gao, Bowen Yu, Chengpeng Li. Dayiheng Liu, Jianhong Tu, Jingren Zhou, Junyang Lin, Keming Lu, Mingfeng Xue, Runji Lin, Tianyu Liu, Xingzhang Ren, and Zhenru Zhang. Qwen2.5-math technical report: Toward mathematical expert model via self-improvement. CORR, abs/2409. 12122, 2024.
[55] Zihao Ye, Lequn Chen, Ruihang Lai, Wuwei Lin, Yineng Zhang, Stephanie Wang, Tianqi Chen, Baris Kasikci, Vinod Grover, Arvind Krishnamurthy, and Luis Ceze. Flashinfer: Efficient and customizable attention engine for LLM inference serving. CORR, abs/2501.01005, 2025.
[56] Jingyang Yuan, Huazuo Gao, Damai Dai, Junyu Luo, Liang Zhao, Zhengyan Zhang, Zhenda Xie, Y. X. Wei, Lean Wang, Zhiping Xiao, Yuqing Wang, Chong Ruan, Ming Zhang, Wenfeng Liang, and Wangding Zeng. Native sparse attention: Hardware-aligned and natively trainable sparse attention. CoRR, abs/2502. 1089, 2025.
[57] Yujia Zhai, Chengquan Jiang, Leyuan Wang, Xiaoying Jia, Shang Zhang, Zizhong Chen, Xin Liu. and Yibo Zhu. Bytetransformer: A perforreforeence transformer boosted for variable- length inputs. In IEEE International Parallel and Distributed Processing Symposium, IPDPS 2023, St. Petersburg, FL, USA, May 15- 19, 2023, pages 344-355, IEBE, 2023.
[58] Lianmin Zheng, Liangsheng Yin, Zhiqiang Xie, Chuyue Sun, Jeff Huang, Cody Hao Yu, Shiyi Cao, Christos Kozyrakis, Ion Stoica, Joseph E. Gonzalez, Clark W. Barrett and Ying Sheng. Sglang: Efficient execution of structured language model programs. In Advances in Neural Information Processing Systems 38: Annual Conference on Neural Information Processing Systems 2024, NeurlPS 2024, Vancouver, BC, Canada. December 10 15, 2024, 2024.
14

---

[59] Size Zheng, Siyuan Chen, Siyuan Gao, Liancheng Jia, Guangyu Sun, Runsheng Wang, and Yun Liang. Tileflow:
framework for modeling fusion dataflow via tree-based analysis. In Proceedings of the 56th Annual IEEF/ACM International Symposium on Microacchitecture, MICRO 2023, Toronto, ON, Canada, 28 October 2023 November 2023, pages 1271-1288.
ACM, 2023.
[60] Yangjie Zhou. Jingwen Leng, Yaoxu Song, Shuwen Lu, Mian Wang, Chao Li, Minyi Guo, Wenting Shen, Yong Li, Wei Lin, Xiangwen Liu, and Hanqing Wu. ugrapher: perfformar graph operator computation via unified abstraction for graph neural networks. In Tor M. Aamodt, Natalie D. Enright Jerger, and Michael M. Swift, editors, Proceedings of the 28th ACM International Conference Architectural Support for Programming Languages and Operating Systems, Volume 2, ASPLOS 2023, Vancouver, BC, Canada, March 25-29, 2023, pages 878- 891. ACM, 2023.
[61] Yangjie Zhou, Wenting Shen, Jingwen Leng, Shuwen Lu, Zihan Liu, Weihao Cui, Zhen- dong Zhang, Wencong Xiao, Baole Ai, Yong Li, Wei Lin, Deze Zeng, Yun Liang, Quan Chen, Ning Liu, and Minyi Guo. Voyager: Input-adaptive algebraic transformations for high- performance graph neural networks. In Proceedings of the 30th ACM International Conference on Architectural Support for Programming Languages and Operating Systems. Volume AS- PLOS '25, page 247-263, New York, NY, USA, 2025, Association for Computing Machinery.
[62] Yangjie Zhou, Honglin Zhu, Qian Qiu, Weihao Cui, Zihan Liu, Peng Chen, Mohamed Wahib, Cong Guo, Siyuan Feng, Jintao Meng, et al. sample-frre compilation framework for efficient dynamic tensor computation. 2025.
[63] Donglin Zhuang, Zhen Zheng, Haojun Xia, Xiafei Qiu, Junjie Bai, Wei Lin, and Shuaiwen Leon Song. Mononn:
Enabling a new monolithic optimization space for neural network inference tasks on modern gpu-centric architectures. In 18th USENIX Symposium on Operating Systems Design and Implementation, OSDI 2024, Santa Clara, CA. USA, July 10-12, 2024, pages 989-1005, USENIX Association, 2024.
15

---

NeurIPS Paper Checklist
1. Claims
Question: Do the main claims made in the abstract and introduction accurately reflect the paper's contributions and scope?
Answer: [Yes]
Justification:
The abstraction and introduction reflect necessary contributions and experi-
ments.
Guidelines:
The answer NA means that the abstract and introduction do not include the claims made the paper.
The abstract and/or introduction should clearly state the claims made, including the contributions made in the paper and important assumptions and limitations. A No or
NA answer to this question will not be perceived well by the reviewers.
The claims made should match theoretical and experimental results, and reflect how much the results can be expected to generalize to other settings.
is fine to include aspirational goals as motivation as long as is clear that these goals are not attained by the paper.
2. Limitations
Question: Does the paper discuss the limitations the work performed by the authors?
Answer: [Yes]
Justification:
The potential limitations is discussed within the paper in separate Sec. 5
Guidelines:
The answer NA means that the paper has no limitation while the answer No means that the paper has limitations, but those are not discussed in the paper.
The authors are encouraged to create a separate "Limitations" section their paper.
The paper should point out any strong assumptions and how robust the results are to violations of these assumptions (e.g.. independence assumptions, noiseless settings, model well-specification, asymptotic approximations only holding locally). The authors should reflect how these assumptions might be violated in practice and what the implications would be.
The authors should reflect on the scope of the claims made, e.g., if the approach was only tested on few datasets or with few runs. In general, empirical results often depend on implicit assumptions, which should be articulated.
The authors should reflect on the factors that influence the performance of the approach. For example, facial recognition algorithm may perform poorly when image resolution is low images are taken in low lighting. Or speech-to-text system might not be used reliably to provide closed captions for online lectures because it fails to handle technical jargon.
The authors should discuss the computational efficiency of the proposed algorithms and how they scale with dataset size.
If applicable, the authors should discuss possible limitations of their approach to address problems of privacy and fairness.
While the authors might fear that complete honesty about limitations might used by reviewers as grounds for rejection, worse outcome might be that reviewers discover limitations that aren't 't acknowledged judgment and recognize that individual actions in favor of transparency play impor- tant role in developing norms that preserve the integrity of the community. Reviewers will be specifically instructed to not penalize honesty concerning limitations.
the paper. The authors should use their best
3. Theory assumptions and proofs
Question: For each theoretical result, does the paper provide the full set of assumptions and complete (and correct) proof?
Answer: [Yes]
16

---

Justification:
The necessary assumption and proof are include in the paper.
Guidelines:
The answer NA means that the paper does not include theoretical results.
All the theorems, formulas, and proofs in the paper should be numbered and cross- referenced.
All assumptions should be clearly stated referenced in the statement of any theorems.
The proofs can either appear in the main paper or the supplemental material, but they appear in the supplemental material, the authors are encouraged to provide short proof sketch provide intuition.
Inversely, any informal proof provided in the core of the paper should be complemented by formal proofs provided in appendix or supplemental material.
Theorems and Lemmas that the proof relies upon should be properly referenced.
4. Experimental result reproducibility
Question: Does the paper fully disclose all the information needed to reproduce the main ex- perimental results of the paper to the extent that affects the main claims and/or conclusions of the paper (regardless whether the code and data are provided or not)?
Answer: [Yes]
Justification:
We have listed detail configuration of experiments in Sec.
Guidelines:
The answer NA means that the paper does not include experiments.
If the paper includes experiments, a No answer to this question will not be perceived well by the reviewers: Making the paper reproducible is important, regardless of whether the code and data are provided or not.
If the contribution is dataset and/or model, the authors should describe the steps taken to make their results reproducible or verifiable.
Depending on the contribution, reproducibility can be accomplished in various ways. For example, if the contribution is a novel architecture, describing the architecture fully might suffice, or if the contribution is a specific model and empirical evaluation, it may be necessary to either make possible for others to replicate the model with the same dataset, or provide access to the model. In general. releasing code and data is often one good way accomplish this, but reproducibility can also be provided via detailed instructions for how to replicate the results, access to a hosted model (e.g.. in the case of large language model), releasing of a model checkpoint, or other means that are appropriate to the research performed.
While NeurIPS does not require releasing code, the conference does require all submis- sions to provide some reasonable avenue for reproducibility, which may depend on the nature of the contribution. For example
(a) If the contribution is primarily a new algorithm, the paper should make clear how to reproduce that algorithm. (b) the contribution is primarily a new model architecture, the paper should describe the architccture clearly and fully.
(c) If the contribution is new model (e.g., a large language model), then there should either way to access this model for reproducing the results or way to reproduce the model (e.g., with an open-source dataset instructions for how to construct the dataset).
(d) We recognize that reproducibility may be tricky in some cases, in which case authors are welcome to describe the particular way they provide for reproducibility. the case of closed-source models, may be that access to the model is limited in some way (e.g., to registered users), but it should be possible for other researchers have some path to reproducing or verifying the results.
5. Open access to data and code
Question: Does the paper provide open access to the data and code, with sufficient instruc- tions to faithfully reproduce the main experimental results, as described in supplemental material?
17

---

Answer: [Yes]
Justification: The code to reproduce the experimental results are provided the supplemen- tary material.
Guidelines:
The answer NA means that paper does not include experiments requiring code.
Please see the NeurlPS code and data submission guidelines (https://nips. cc/
for more details.
While we encourage the release of code and data, we understand that this might not be possible, so "No" is an acceptable answer. Papers cannot be rejected simply for not including code, unless this is central to the contribution (e.g., for new open-source benchmark).
The instructions should contain the exact command and environment needed to to reproduce the results. See the NeurIPS code and data submission guidelines (https
httpolliliggideeeCoooooo /ooeeee/e/ooeubbbiiieoo ssicnlolicy) for more details.
The authors should provide instructions on data access and preparation, including how to access the raw data, preprocessed data, intermediate data, and generated data, etc.
The authors should provide scripts to reproduce all experimental results for the new proposed method and baselines. only subset of experiments are reproducible, they should state which ones are omitted from the script and why.
At submission time, preserve anonymity, the authors should release anonymized versions (if applicable).
Providing as much information as possible in supplemental material (appended to the paper) recommended, but including URLs to data and code permitted.
6. Experimental setting/details
Question: Does the paper specify all the training and test details (e.g., data splits, hyper- parameters, how they were chosen, type of optimizer, etc.) necessary to understand the results?
Answer: [Yes]
Justification: This not the core contribution of this work. But still, we follow the standard method of prior work as we have mentioned in the Sec. 4
Guidelines:
The answer NA means that the paper does not include experiments.
The experimental setting should be presented in the core the paper to a level of detail that is necessary to appreciate the results and make sense of them.
The full details can be provided either with the code, appendix, as supplemental material.
7. Experiment statistical significance
Does the paper report error bars suitably and correctly defined or other appropriate information about the statistical significance of the experiments?
Question:
Answer:
[Yes]
Justification: The experiments are conducted many times and report average results,
Guidelines:
The answer NA means that the paper does not include experiments.
The authors should answer "Yes" if the results are accompanied by error bars, confi- dence intervals, or statistical significance tests, at least for the experiments that support the main claims of the paper.
The factors of variability that the error bars capturing should be clearly stated (for example, train/test split, initialization, random drawing of some parameter, or overall with given experimental conditions).
The method for calculating the error bars should be explained (closed form formula, call to a library function, bootstrap, etc.)
The assumptions made should be given (e.g., Normally distributed errors).
18

---

It should be clear whether the error bar is the standard deviation or the standard error of the mean.
is OK to report 1-sigma error bars, but one should state it. The authors should preferably report a 2-sigma error bar than state that they have 96% CI, if the hypothesis of Normality of errors is not verified.
For asymmetric distributions, the authors should be careful not to show in tables or figures symmetric error bars that would yield results that are out range (e.g. negative enror rates).
If error bars reported in tables plots, The authors should explain in the text how they were calculated and reference the corresponding figures or tables in the text.
8. Experiments compute resources
For each experiment, does the paper provide sufficient information on the com- puter resources (type of compute workers, memory, time of execution) needed reproduce the experiments?
Question:
Answer: [Yes]
Justification:
The specific configurations are included in the Sec.
Guidelines:
The answer NA means that the paper does not include experiments.
The paper should indicate the type of compute workers CPU or GPU, internal cluster, or cloud provider, including relevant memory and storage.
The paper should provide the amount compute required for each of the individual experimental runs well estimate the total compute.
The paper should disclose whether the full research project required more compute than the experiments reported in the paper (e.g., preliminary failed experiments that didn't make into the paper).
9. Code ethics
Question: Does the research conducted the paper conform, in every respect, with the NeurlPS Code of Ethics http:: httppp////eerrr....co//////////a
Answer: [Yes]
Justification:
We follow the NeurIPS Code of Ethics.
Guidelines:
The answer NA means that the authors have not reviewed the NeurIPS Code of Ethics.
If the authors answer No, they should explain the special cireumstances that require deviation from the Code of Ethics.
The authors should make sure to preserve anonymity (e.g., if there is special consid- eration due to laws or regulations their jurisdiction..
10. Broader impacts
Question: Does the paper discuss both potential positive societal impacts and negative societal impacts of the work performed?
Answer: [Yes]
Justification: We discuss the potential positive societal impacts and negative societal impacts the work in the paper.
Guidelines:
The answer NA means that there no societal impact of the work performed.
If the authors answer NA or No, they should explain why their work has no societal impact why the paper does not address societal impact.
Examples of negative societal impacts include potential malicious or unintended uses (e.g., disinformation, generating fake profiles, surveillance), fairness considerations (e.g., deployment of technologies that could make decisions that unfairly impact specific groups), privacy considerations, and security considerations.
19

---

The conference expects that many papers will be foundational research and not tied to particular applications, let alone deployments. However, if there is direct path to any negative applications, the authors should point it out. For example, it is legitimate to point out that an improvement in the quality of generative models could be used to generate deepfakes for disinformation. On the other hand, it is not needed to point out that a generic algorithm for optimizing neural networks could enable people to train models that generate Deepfakes faster.
The authors should consider possible harms that could arise when the technology is being used as intended and functioning correctly, harms that could arise when the technology being used as intended but gives incorrect results, and harms following from (intentional or unintentional) misuse of the technology.
If there negative societal impacts, the authors could also discuss possible mitigation strategies (e.g., gated release of models, providing defenses in addition to attacks, mechanisms for monitoring misuse, mechanisms monitor how system learns from feedback over time, improving the efficiency and accessibility of MI..
11. Safeguards
Does the paper describe safeguards that have been put in place for responsible release of data models that have high risk for misuse (e.g., pretrained language models, image generators, or scraped datasets)?
Question:
Answer: [NA]
Justification:
The answer NA means that the paper poses no such risks.
Guidelines:
The answer NA means that the paper poses such risks.
Released models that have high risk for misuse or dual-use should be released with necessary safeguards to allow for controlled use of the model, for example by requiring that users adhere usage guidelines or restrictions to access the model or implementing safety filters.
Datasets that have been scraped from the Internet could pose safety risks. The authors should describe how they avoided releasing unsafe images.
We recognize that providing effective safeguards is challenging, and many papers do not require this, but we encourage authors to take this into account and make best faith effort.
12. Lieenses for existing assets
Are the creators or original owners of assets (e.g., code, data, models), used in the paper. properly credited and are the license and terms of use explicitly mentioned and properly respected?
Question:
Answer: [Yes]
Justification: We cite the original paper that produced the code package dataset.
Guidelines:
The answer NA means that the paper does existing assets.
The authors should cite the original paper that produced the code package or dataset.
The authors should state which version of the asset is used and, if possible, include URL.
The name of the license (c.g., CC-BY 4.0) should be included for each asset.
For scraped data from a particular source (c.g., website), the copyright and terms of service of that source should be provided.
If assets are released, the license, copyright information, and terms of use in the package should be provided. For popular datasets, paperswi thcode. com com/datasets has curated licenses for some datasets. Their licensing guide can help determine the license of dataset.
For existing datasets that are re-packaged, both the original license and the license of the derived asset (if it has changed) should be provided.
20

---

If this information is not available online, the authors are encouraged to reach out to the asset's creators.
13. New assets
Question: Are new assets introduced in the paper well documented and is the documentation provided alongside the assets?
Answer: [NA]
Justification: This paper does not release new assets.
Guidelines:
The answer NA means that the paper does not release new assets.
Researchers should communicate the details of the dataset/code/mmoel as part of their submissions via structured templates. This includes details about training, license, limitations, etc.
The paper should discuss whether and how consent obtained from people whose asset is used.
At submission time, remember to anonymize your assets (if applicable). You can either create an anonymized URL or include an anonymized zip file.
14. Crowdsoureing and research with human subjects
Question: For crowdsourcing experiments and research with human subjects, does the paper include the full text of instructions given to participants and screenshots, if applicable, as well details about compensation (if any)?
Answer: [NA]
Justification: This paper does not involve crowdsourcing nor research with human subjects.
Guidelines:
The answer NA means that the paper does not involve crowdsourcing nor research with human subjects.
Including this information in the supplemental material is fine, but if the main contribu- tion of the paper involves human subjects, then as much detail as possible should be included the main paper.
According the NeurIPS Code of Ethics, workers involved in data collection, curation, or other labor should be paid at least the minimum wage in the country of the data collector.
15. Institutional review board (IRB) approvals or equivalent for research with human subjects
Question:
Does the paper describe potential risks incurred by study participants, whether such risks were disclosed to the subjects, and whether Institutional Review Board (IRB) approvals (or an equivalent approval/review based on the requirements of your country or institution) were obtained?
Answer: [NA]
Justification: This paper does not involve crowdsourcing nor research with human subjects.
Guidelines:
The answer NA means that the paper does not involve crowdsourcing nor research with human subjects.
Depending the country in which research is conducted, IRB approval (or equivalent) may be required for any human subjects research. If you obtained IRB approval, you should clearly state this in the paper.
We recognize that the procedures for this may vary significantly between institutions and locations, and expect authors to adhere to the NeurlPS Code of Ethics and the guidelines for their institution.
For initial submissions, do not include any information that would break anonymity (if applicable), such the institution conducting the review.
16. Declaration of LLM usage
21

---

Question: Does the paper describe the usage of LLMs if it is important, original, or non-standard component of the core methods this research? Note that if the LLM used only for writing, editing, or formatting purposes and does not impact the core methodology, scientific rigorousness, or originality of the research, declaration is not required.
Answer: [NA]
Justification: LLM is used only for writing and editing in this paper.
Guidelines:
The answer NA means that the core method development in this research does not involve LLMs any important, original, non-standard components.
Please refer to our LLM policy fattps://neurd/: /ttpr//pp...co/oofeeeeee//////// for what should or should not be described.
22

---

A Related Works
Operator-Level Optimizations
Numerous researchs have explored operator-level optimizaitons
61 62 FlashAttention / fuses the entire Attention operation into a single memory kernel. Building on this, FlashDecoding [10] extends parallelism to the KV cache sequence dimension during decoding. FlashDecoding++ [9] proposes to determine the scaling factor based on statistics in advance and introduces FlatGEMM to represent the GEMM with highly reduced dimension in decoding. FlashAttention: [44] demonstrated that warp specialization and new hardware features [33] such as asynchrony can have a significant impact on the Attention.
for LLM inference [60) 30 28 21
FLashMLA [22] design an efficient decoding Attention kernel for DeepSeek MIA architecture inspired by FlashAttention-3. However, there optimizaitons adopt a block-isolated execution pattern. Due to the lack of structured communication across thread blocks, intermediate results are repeatedly written to and read from global memory, limiting opportunities for broader operator fusion and on- chip reuse. Cl ClusterFusion explores a more general fusion space enabled by efficient cluster-scoped collective primitives.
Algorithm-Level Optimizations Several studies focus on improving LLM inference efficiency through algorithmic changes [26] 29 56 31]. Techniques such as quantization and sparsification aim to reduce computational and memory overhead. Quantization 266 29] compresses model weights and activations by converting high-bitwidth representations into lower-bitwidth ones, reducing arithmetic and memory costs. Pruning [56] 31 increases the proportion of zero elements in weights or activations, enabling hardware to skip redundant computations. These techniques are orthogonal to work. ClusterFus in does not modify the model workload but can complement these approaches by optimizing the underlying dataflow and kernel execution through structure and reusable primitives.
Systems on Inter-Core Connected Hardwares The on-chip inter-core interconnect has also been adopted in alternative hardware platforms such as Graphcore IPU [20] and Cerebras WSE (6]. Several prior works have explored leveraging this capabilities to optimize deep learning workloads. T10 [27 an end-to-end deep learning compiler targeting inter-core connected intelligence processors. with an emphasis on fine-grained modeling of data movement between cores. WaferLLM (18), on the other hand, is the first system to propose a I.LM parallelism solution tailored for wafer- scale accelerators. Moreover, these systems are closely tied to custom hardware architectures that offer full inter-core connectivity or integrates inter-core communication mechanisms into specific operators like GEMM and GEMV. As such, their design assumptions do not translate well to modern GPU architectures like NVIDIA Hopper which remains the dominant platform for LLM inference deployment. ClusterFusion on introduces cluster-scoped communication primitives which abstract common aggregation and reduction patterns and can be flexibly reused across different workloads.
B Additional Dataflows
In this section, we first introduce the computation process of DeepSeck MLA and present our fused MLA dataflow, which leverages cluster-level collective communication primitives, evaluated in the main paper. We then describe an alternative dataflow in which thread blocks partition the head dimension within the Attention module, in contrast to the dataflow in the main paper that partitions the KV cache along the token dimension, demonstrating the capability of our primitives in enabling different dataflows for the same computation. Finally, we evaluate these dataflow variants by analyzing their corresponding DSMEM traffic.
B.1 DeepSeek MLA
DeepSeek introduces the Multi-Head Latent Attention (MLA) mechanism, which has been adopted in its model series [1] [2]. During inference, weight absorption optimization 43 is applied in the decoding stage of MLA to reduce overall computational cost. The detailed computation processes of both the original MLA and the optimized version in DeepSeek- BeepSeek-V2-Liter 11] are illustrated in Fig. [4] In the original MLA computation, the input hidden state first undergoes a Q Projection to generate the multi-head Q and Down Projection to obtain the compressed KV cache. The newly generated compressed KV then concatenated with the cached KV and passed through an Up Projection to produce the multi-head head KV, which is used in the MHA module. In the weight absorption version of
23

---

MLA Basic
MLA Weight Absorption
CHALPP
1.912
T-1
1:16 912
1, 11..1.00-000
T
5, 512
T+ 1.0
Multi-Head Attention (MHA)
Multi- Multi-Query Attention (MQA)
1,
Figure 14: Overview of MLA computation: original (left) and with weight absorption optimization (right). The black stipple represents the newly generated latent presentation, which will be cached to calculate keys and values for Attention computation.
Hiedon
182
185
Ouepue
Weight
Weight boundary Colerice.synCD
Weigt leducs
GeMV
Softmax
Head
Head
Onnohio
Gothr
Figure 15: Cluster-centric fused MLA dataflow graph.
MLA, the process begins by computing and K, where V is obtained by partially reusing K:
Q = Hidden x Wq x Wup K = Hidden x WK, V K[: kv llkllola_rraaaaalalllllallaaaaee rank]
(3)
Next, the Atlention module is computed as:
(QIKT V... V....,
QQK Vilx
Z = Concat
Softmax
Softmax
v
(4)
ViTy
Finally, the result is passed through Down Projection to produce the final Attention output:
Output = Z x Down
(5)
The key difference between MLA and conventional MHA lies in the introduction of additional projection layers, specifically the Up Projection Down Projection, which are designed to preserve mathematical consistency. Additionally, MIAA employs the MQA mechanism (46. in which all Q heads share a single KV cache head. This approach increases computational intensity while reducing the memory access associated with the KV cache. Moreover, the head dimension in MLA corresponds to the value of kv aollllrrrraa. which is typically larger than that in other models. For example, in DeepSeck-V2-Lite. it is 512, whereas in Llama2-7B, the head dimension is 128.
According to our cluster-centri dataflow design principle, the fused MLA dataflow is illustrated in Fig. [5] This dataflow performs the entire MLA computation on-chip, eliminating any intermediate off-chip memory traffic. The detailed algorithm is presented in Alg. This dataflow parallelized across attention heads, with each head assigned to a cluster consisting of N 2 thread blocks
(k: < 4). Each thread block within a cluster is assigned rank b [0, - 1]. Within each cluster, thread blocks collaboratively partition the head dimension for the QKV Projection; the

---

Algorithm 4 Fused MLA Dataflow -Thread Block View
Require: Input hidden states H, RIXD, Q Projection weight W E Dx KV Projection weight
W E Dx Up Projection weight w E RHxx, Down Projection weight W, Down
R Output Projection weight W E RHxd and KV cache KV R in global memory.
1: Allocate shared memory buffers: Qs, KV, E Rixi and Sum, Smax (sofumax statistics).
2: Compute segment results of Q Projection: Qs H W
3: Compute segment results of KV Projection: KV, H. x W
4: Obtain the complete QKV: Q + ClusterGather Qo), KV, + ClusterGather (KV)))
5: Compute segment results of Up Projection by batch matmul: Qs + Q W
6: Obtain the complete Q: Q2 + ClusterGather (Qa).
7: Compute partial result of Attention similar the FlashDecoding dataflow:
Compute S, f exp( Q (KVa KVb)"), obtain local Saum Smax.
And store Smax in register Reg
> softmax statistics.
Compute A - S₁ x (KVg KV&) b As reuse the shared memory space of Qs
8: Obtain the complete softmax statistics Sum and Smax:
Saum ClusterReduce sum), Smax <- www.sleeeeeeuueSSSnnn.. max)
9: Rescale Attention output according to online softmax:
As AG exp( Reg Smax)
Saum
Obtain the complete Attention output: A + ClusterReduce (A,, sum).
11: Compute partial results of Down Projection by batch matmul: A& f- A, x WPonn
Obtain the complete Down Projection output: As + ClusterReduce(AA (A,, sum).
13: Compute the result of the Output Projection and write it to global memory using atomicAdd:
O₂ + Asx
14: return Og.
kv a_lloor_raann rank dimension for the Up Projection and Down Projection modules; the token dimension of the KV cache for the Altention module; and the output dimension for the Oulput Projection. And each thread block processes the full input hidden states and computes its corresponding output tile On after the Output Projection. In this algorithm, B denotes the batch size, D the input hidden dimension, and H the total head dimension. The variables h, I, s, and d refer to the partitioned sizes of the head dimension, kv kv_lora_raal rank, sequence length, and output dimension per thread block, respectively. For simplicity, we omit the rope_dim shown in Fig. 14
We also estimate the total DSMEM traffic incurred by the dataflow used in Alg. 4 which involves three ClusterGather operations and three ClusterReduce operations. The DSMEM traffic for the ClusterGather operations is given by:
Traf fiCGather (h, N) + 2 x ficcthecch
and the traffic for the ClusterReduce operations is:
Traf FrafffccRede fic (1, NN) + frof f FicReduce (H,
We omit the traffic introduced by the softmax statistics, as involves only two floating-poin values and is negligible compared to the tensor-level data movement.
B.2 SplitHead Dataflow
In line with our cluster-centric dataflow design principles and by leveraging cluster-level communica- tion primitives, we implement the fused QKV Projection, Attention, and Output Projection dataflow described the main paper, where intermediate data is stored on-chip shared memory. addition, we design alternative dataflow called SplitHead dataflow that stores the intermediate data in faster
25

---

Algorithm 5 SplieHead Dataflow -Theedd Block View
Require: Input hidden states Hg E BRBDD QKV Projection weight W E RD Outpur Projection weight W E R and KV cache Kg V E R in global memory.
1: Allocate register memory buffers: Ko Vo E Rxhh shared memory buffer: S, E R
2: Compute segment results of QKV Projection: Qs, Kb, Vo t H, x W W
3: Compute segment result of Altention:
Compute S, -- Q x (Kg K²)²
Obtain the complete result of Q x KT: Sg ClusterReduce sum).
Compute Softmar: S, &- softmazt (Sp).
Compute Attention output: As + S₁ x (Vg Vv)
4: Compute the partial one head result of the Outpul Projection: O, + A. x Wo.
5: Obtain the complete one head result of the Outpul Projection:
O ClusterReduce(O,, sum)
6: Write the complete result of the Output Projection to global memory using atomi cAdd:
atdmi
f atomicAdd( O)
7: return O.
Hidden
Fused QKV_Pro) Proj_Attn_Ouu Proj
Cluster Reduce
TB0
TB1
TB2
Output
TB3
TB4
TB5
TB6
Partial inner-product
Weight Q/K/V
Key Cache
lveq_een, 1)
Value Cache
Weight
GeMV
Softmax
Rescale
Head
Head
Head 3 On-chip
Kernel boundary
(device. sonice..ync))
Figure 16: Cluster-centric fused QKV Projection, Attention and Output Projection dataflow graph.
on-chip register memory. As illustrated in Alg. 5 this dataflow is parallel in the number of heads.
Each head corresponds to cluster of N = 2* thread blocks (I < 4), where each block is assigned rank b E N0, N - 1]. Within each cluster, thread blocks just partition the head dimension in QKV Projection, Attention, and Outpul Projection. For the whole dataflow, each thread block processes the entire input hidden states and computes the corresponding partial output O, after the Output Projection. In this algorithm, B denotes the batch size, / the input hidden dimension, H the total head dimension, S the KV cache sequence length, and h represent the partitioned sizes of the head dimension and the output dimension, per thread block, respectively. In this design, need to reduce the result of Q x K" which has a shape of Sequence Length x Batch Size. Each thread block holds the segment Attention output and computes the partial Output Projection, which must then be reduced and written to global memory using atomi cAdd.
As shown in Fig. 16, this dataflow enables intermediate results such as Qo. K,, and V. to be stored in faster register memory, instead of shared memory. However, the total DSMEM traffic incurred by this dataflow is
Traf Prafficraat = Traf I fiCneduce (S, N) + DrraffTarcaaaaaa SiCneducc (D, N)
which higher than the dataflow that partitions the KV cache along the token dimension in the Attention module, proposed in the main paper and Sec. B.I The total DSMEM traffic of the
26

---

sGLans VLLM
TRI-LLM
SCLang MLCC-TM
VLLM
TRE-LLM
ClusierFusion
ClusterFusion
2000
800
TPOT (ms)
60-
Latency (us)
1500
40-
1000
400
500
200
0
16k
2k
Sequence
Length
Sequence Length
Sequence Length
Sequence Length
Figure 17: TPOT of Llama2-7B (left) and Figure 18: Latency of core modules of Llama2- 7B (left) and DeepSeek V2-Lite (right) on H100.
DeepSeck- CrepSeek-V2-Litte V2-Lite (right) on H100.
Kammal Lauadd Ovatheal (mns)
Momary Thanstor Sias (MBB)
SpirTckon
Elama--0 -88
Duyplak- Vil
6.0
vLLM
SpritHeed
SOilann
CusteFusia
Latency (us)
200
Liama
DeepSook
300
600
1.04
1550
200
2000
100
DeugSook
1k
8k
16K
1K
16k Sequenoe Longth
4k
8k
Soquence Length
Sequende Length
Models
Sequence Longth
Figure 19: Comparison of global memory data transfer size
Figure 20: Latency comparision of SplitToken and SplitHead dataflows.
(left) and GPU kernel launch overhead (right).
SplitToken dataflow is
Traf fictotal T f ficReduce (H, N) Traf ficGather (3h, .N)
which is significantly lower, mainly depends on the head dimension H the kv kv_lorss v__lora_raan, in the fused MLA dataflow, both of which are much smaller than the sequence length S that dominates the DSMEM traffic in the SplitHead dataflow. The increased communication overhead outweighs the benefits of using register memory. As demonstrated by our experimental results presented later, the SplitHead dataflow yields higher latency.
C Additional Experiments
C.1
Multi-Batch Evaluation and Speedup Analysis
We conduct additional experiments using a batch size of 16, while keeping other experimental settings consistent with those described in the main paper. The TPOT and latency results for core modules are shown in Fig. 17 and Fig. [8] respectively. On average, ClusterFusic achieves speedups of 1.11x, 1.09x, 1.12x, and 1.32x over SGLang, vLLM, TensorRT-LLM, and MLC-LLM. respectively, across various sequence lengths on Llama2-7B with a cluster size of 4. For DeepSeek- V2-Lite, ClusterFusio delivers average speedups of 1.15x, 1.14x, 1.07x, and 1.84x under the same conditions. In terms of the core modules, including fused QKV Projection, Attention, and Output Projection, ClusterFusion achieves average speedups of 1.14x, 1.12x, 1.2x, and 1.41x over SGLang, vLLM, TensorRT-LLM, and MLC-LLM, respectively, on Llama2-7B, Similarly, for DeepSeek- DeepSeek-V2-Lite, the corresponding speedups are 1.19x, 1.18x, 1.14x, and 2.04x.
We identify two primary factors contributing to the performance advantage of ClusterFusion over existing frameworks with CUDA Graph optimizations: reduced global memory transfer volume and significantly lower GPU kernel launch overhead 63 [53). To quantify these henefits, we use NVIDIA Nsight Systems (38] and Nsight Compute [37] to profile memory traffic and kernel launch overhead under batch size of 16. The results are presented in Fig, 19 The observed performance gains are primarily attributed to ClusterFusio executing QKV Projection, Attention, and Output Projection entirely on-chip, thereby mi nimizing intermediate memory traffic. Furthermore, ClusterFusion reduces kernel launch overhead by nearly an order of magnitude in end-to-end scenarios, even compared to baselines already optimized with CUDA Graph. However, the reduction in global memory traffic has limited impact the multi-batch scenario, as the KV cache and model weights still dominate memory usage, while the intermediate memory footprint remains small. Moreover, the overall computation intensity increases significantly with larger batch sizes, leading to a reduced speedup compared to the single-batch results presented in the main paper.
27

---

C.2 SplitHead Dataflow Evaluation
We also implement the SplitHead dataflow described in Sec. B.2 and present the performance comparison between the SplitToken, SplitHead dataflows and two representive baselines in Fig. [20] When the sequence length is short, the latency difference is minimal because intermediate data can be stored in register memory, which improves efficiency compared to the SplitToken, and the gap DSMEM traffic compared to the SplitToken dataflow remains small. However, as the sequence length increases, the DSMEM traffic of the SplitHead dataflow grows significantly larger than that of SplitToken, resulting in increased latency. From the perspective of operator fusion efficiency, the SplitHead dataflow enables fusion with intermediate data stored in high-speed register memory. as long as the data size is small enough to entirely within the registers. However, despite this benefit, the SplitHead dataflow incurs significantly higher DSMEM traffic, especially as the sequence length increases. This increased traffic becomes a major bottleneck and leads to worse overall performance. Therefore, our cluster-centri dataflow design takes into account not only fusion efficiency but also the memory traffic different dataflows, aiming to achieve better overall performance.
28


Reviews:
[{"review_id": "E3m9WDlVRQ", "reviewer": "Reviewer_q92t", "summary": "The paper addresses decoding latency in LLM inference caused by fragmented kernels and frequent off-chip memory sync. Although NVIDIA Hopper GPUs offer low-level DSMEM, they lack high-level communication primitives.  \n\n**ClusterFusion** fills this gap by introducing two on-chip collective operations:  \n- **ClusterReduce**  \n- **ClusterGather**  \n\nThese enable QKV, attention, and output projections to remain on-chip, cutting off-chip transfers and boosting performance. On H100 GPUs with Llama2-7B and DeepSeek-V2-Lite, ClusterFusion achieves a **1.61×** speedup over SGLang, vLLM, TensorRT-LLM, and MLC-LLM.\n\nThe paper contributes:\n1. **Profiling** of decoding dataflow inefficiencies.  \n2. **High-level primitives** for structured inter-block communication on Hopper.  \n3. **ClusterFusion framework** that expands operator fusion and reduces inference latency.", "questions": "1. **Cluster Size (k) Selection**  \n   - **Issue:** The paper lacks concrete guidance on choosing the cluster size \\(k\\).  \n   - **Request:** Provide an analytical or discussion for \\(k\\) selection and for figure 11, could you deliver a comprehensive for more different models (e.g., reasoning model)\n\n2. **Experiment Clarify**\n   - **Issue:** It's unclear for FlashInfer [45], Triton [40] and TVM usage.\n   - **Request:** Could you clearly tell me for vLLM, SGlang, etc, what's the actual kernel we are using for each framework. For example, do we use triton-based kernel or we do use torch.compile generated kernel?\n\n3. **FFN analysis**\n   - **Issue** I am confused why author also use these two primitives for MLP.\n   - **Request** Could you give ma an analysis or discussion for using these two primitives for MLP? And what's the potential benefits?\n\n4. **Evaluation on Long Contexts and Scalability to More Models**  \n   - **Issue:** Claims of benefits for long-sequence decoding are not validated on truly large contexts and larger / multi-modality model.  \n   - **Request:** Report latency and accuracy on models handling > 100 K token contexts (e.g., DeepSeek R1, Qwen-Long, Llama-4) to substantiate your scalability claims.  \n    - **Request:** Provide empirical results or detailed performance projections for multi modality or larger model (e.g., llama 3.3 70 B), and discuss any emerging bottlenecks or diminishing returns.  \n   \n5. **Cross-Architecture Portability**  \n   - **Issue:** The approach is demonstrated only on NVIDIA Hopper GPUs.  \n   - **Request:** Discuss architectural assumptions and outline how ClusterReduce/ClusterGather could be implemented on other platforms (e.g., AMD/Intel GPUs or TPUs). If possible, include a small-scale result on one alternative accelerator.", "limitations": "Yes. There is no potential negative societal impact of this work.", "rating": 4, "confidence": 4, "soundness": null, "presentation": null, "contribution": null, "strengths": "", "weaknesses": "", "quality": 4, "clarity": 4, "significance": 2, "originality": 2, "strengths_and_weaknesses": "## Strengths\n\n1. **Quality**  \n   - The paper delivers a rigorous motivation analysis to algorithm design. I do love the diagrams, pseudocode and smooth storyline of paper that walk me through on-chip end2end design.\n   - Empirical measurements are comprehensive—covering per-kernel breakdowns, DSMEM bandwidth vs. global memory, and SM-to-SM link utilization.\n\n2. **Clarity**  \n   - **ClusterReduce** and **ClusterGather** are crisply specified: their APIs map directly to common data-parallel patterns (all-reduce, broadcast) and hide DSMEM complexities.  \n\n3. **Significance of Results**  \n   - The approach first proposes such primitive and directly targets production-critical workloads. On real hardware (NVIDIA H100), ClusterFusion yields a consistent **1.61×** end-to-end speedup over four competing inference frameworks.\n\n---\n\n## Weaknesses\n\n1. **Originality vs. Engineering Effort**  \n   - The primitives rest on NVIDIA PTX’s DSMEM features; while the abstraction is valuable, it does not introduce a fundamentally new communication algorithm or theoretical model. The novelty is primarily in the high-level API and scheduler integration rather than in core algorithms.\n\n2. **Generality Across Architectures**  \n   - The paper mentioned different architecture may suffer from these issues, but all experiments target Hopper GPUs. It remains unclear how ClusterReduce/Gather would perform—or even map—to architectures without hardware DSMEM (e.g., AMD ROCm, TPUs). A discussion or prototype on alternate hardware would bolster claims of broad applicability.\n\n3. **Scalability to Very Large Models and Contexts**  \n   - Evaluation caps at 7 B-parameter models and 16 K token contexts. It is left open whether on-chip buffers suffice for, say, 70 B models or 100 K-token sequences without falling back to global memory, which would erode benefits.\n\n4. **Depth of Algorithmic Detail**  \n   - The paper’s pseudocode for collective setup omits corner-case handling (e.g., SM load imbalance, DSMEM bank conflicts). A deeper dive into these implementation subtleties would aid reproducibility.  \n   - Latency models for DSMEM vs. global memory are presented qualitatively; a formal performance model (e.g., roofline-style) could clarify when the method breaks even.\n\n5. **Clarity in Presentation of Limitations**  \n   - While the authors briefly note DSMEM capacity limits, they do not quantify how often fusion must be “split” across multiple passes or how that impacts latency.", "was_revised": true, "final_justification": "The technical concerns I had with this paper are mostly resolved. I will raise my point.\n\nBut I hope to see the final manuscript include our discussion in the full paper/appendix and include the points we discussed during rebuttal.", "rebuttal": "Thank you for your thoughtful feedback on our paper. We would like to provide more clarification.\n\nTo answer your concerns:\n1. **Cluster Size (k) Selection.** Based on our profiling results in Section 2.3 and Figure 5, we observe that using an excessively large cluster size can degrade performance due to reduced effective SM bandwidth and lower active SM. To mitigate this, we recommend choosing a cluster size less than or equal to 8 for most workloads. Additionally, to facilitate data tiling and ensure alignment between data layout and thread-block mapping, we suggest using even cluster sizes such as 2, 4, or 8—resulting in a small but practical space. We also expose cluster_size as a tunable parameter in our implementation, since its optimal value is tightly coupled with several factors such as tile size and kernel-level parallelism. In practice, this introduces a trade-off: when the cluster count is large and each cluster handles a small amount of data, a smaller cluster size leads to sufficient parallelism and more favorable tiling. On the other hand, when each cluster’s workload is relatively large, a larger cluster size is needed to efficiently amortize inter-block communication. This behavior is reflected in our results in Section 4.1 and Figure 11, where smaller cluster sizes (2 or 4) are preferred for many-head attention, while larger ones become more favorable as the head count decreases or head dimension grows. \n2. **Kernels used in each framework.** In our experiments, we adopt the default operators used by each framework under their officially recommended configuration[4,5,6,7], with `torch.compile` enabled to utilize optimized Triton-based kernels for decoding. The following table summarizes the actual attention and GEMM kernels used in each baseline framework:\n\n| Framework      | Attention Kernel                                    | QKV/FFN GEMM Kernel                    |\n|----------------|-----------------------------------------------------|----------------------------------------|\n| **SGLang**     | FlashInfer (`BatchDecodeWithPagedKVCacheWrapper`)  | Torch.compile generated Triton Kernel |\n| **vLLM**       | FlashAttention-3 (`flash_attn_with_kvcache`)        | Torch.compile generated Triton Kernel|\n| **TensorRT-LLM** | FlashAttention-3 (`flash_attn_with_kvcache`)      | Torch.compile generated Triton Kernel  |\n| **MLC-LLM**    | TVM-generated kernel                                | TVM-generated kernel                  |\n\nAll of these kernels represent SOTA implementations in their respective frameworks. Our method surpasses their performance, demonstrating both the generality and effectiveness of our approach. We will include detailed information about these kernel in the final version of the paper.\n\n3. **About FFN.** We have extended our primitives to fuse other modules like FFN, and the corresponding implementation(codes) is included in our supplementary material. In short, FFN fusion is feasible, but the performance gain is limited. \nFirstly, the FFN kernel must be separated from the Output projection, because the output projection in Attention introduces a all-to-one data dependency: all thread blocks must reduce their output to the same address. This is not feasible to resolve fully on-chip via DSMEM without falling back to global memory, as the limited cluster size. Secondly, we designed a fused FFN kernel where we partition the intermediate size across blocks within a cluster, and split the output dimension across different clusters. ClusterReduce is then used to compute partial sums across blocks on chip, followed by local element-wise SiLU and residual addition. However, due to the large intermediate dimensions in most mainstream LLMs (e.g., 11008 or higher), each thread block must handle a huge reduction payload. As shown in our profiling results (Section 2.3 and Figure 5), the limited DSMEM bandwidth in this case prevents us from gaining significant on-chip communication benefits. However, FFN fusion still helps reduce kernel launch overhead and global memory traffic and achieves comparable latency to existing frameworks:\n\n| Framework         | FFN Latency (us) |\n|-------------------|------------------|\n| SGLang            | 115  |\n| vLLM              | 113|\n| TensorRT-LLM      | 107 |\n| MLC-LLM           | 110|\n| Ours  | **118** |\n\n4. **Evaluation on longer context and larger models.** First, regarding long-context capability, we have extended our evaluation of Llama2-7B to >100K context length. In this setting, the TPOT(ms) of our fused kernel against baselines is as follows:\n \n| Seqlen\\Framework | SGLang | vLLM | Ours |\n| ------ | ------ | ------ | ------ |\n| 128k   | 33.5   | 33.7 | **25.6** |\n\nNow, we have already evaluated a range of model sizes under common single-GPU inference settings. As for scalability to larger and multi-modal models with long context that require distributed inference, we have confirmed that our kernel maintain measurable per-device speedups even when integrated into a multi-GPU inference pipeline. As an initial step, we tested a distributed setup for DeepSeek-R1-671B with TP=16, where our single-GPU fusion kernels still yield measurable gains in the per-device computation stage. In this setting, the TPOT(ms) of our fused kernel against baselines is as follows:\n\n| Seqlen\\Framework | SGLang | vLLM | Ours     |\n| ------ | ------ | ------ | ------ |\n| 128k   | 38.1   | 38.5 | **37.1** |\n\nWhile these results are currently limited to per-device computation time, they demonstrate the effectiveness of our approach for integration into multi-GPU/node inference. When scaling to larger models such as DeepSeek-R1, the speedup diminishes slightly (e.g., from 1.3× on LLaMA2-7B to ~1.1× on R1). This is because operator fusion has less room to eliminate overheads in larger models. Our future work focus on the extensions that co-optimize communication and memory access which will further improve performance. To address cross-device scalability, our implementation allows seamless insertion of collective communication operations (e.g., ring-based all-reduce, gather) without disrupting kernel scheduling. This enables future support for end-to-end fusion across both intra- and inter-GPU workloads with minimal integration effort.\n\n5. **Generality on other hardwares.** The techniques and rationale implemented in ClusterFusion are not limited to NVIDIA Hopper GPUs. Our fundamental goal is to enable holistic dataflow design at the larger scale. As discussed in the background section, existing fragmented design makes it difficult to achieve globally optimal. We construct ClusterReduce and ClusterGather that enable on-chip cooperation between thread blocks. The primitives can be mapped to other platforms. On AMD GPUs, SMs, shared memory, and DSMEM used in our algorithm can be naturally mapped to compute units, local data share (LDS), and LDS-to-LDS interconnects (e.g., in RDNA 4[1]). On dataflow architectures such as Graphcore IPUs [2] and Cerebras Wafer-Scale Engines [3], our fine-grained task partitioning strategy remains applicable across the many-core compute fabric of these systems. The holistic dataflow design philosophy as well as our inter-core tree-pattern ClusterReduce and ClusterGather can be further explored on these platforms to fully exploit their on-chip interconnects and distributed memory hierarchy. We are actively collaborating with engineers to implement our primitives on other GPU like MI300X/MI308X GPUs, enabling the same holistic, end-to-end dataflow optimization.\n6. **About novelty.** While reduce and gather are common primitives in distributed system, they are rarely considered in chip-scale with interconnect. This enable new optimization scope and opportunities, and our novelty mainly lies in the exploration of on-chip dataflow via interconnect. On Hopper GPUs, blocks within a cluster are fully connected via DSMEM, as illustrated in Figure 4 and Figure 6. And as shown in Section 2.3 and Figure 5, DSMEM performance is non-trivial and not automatically efficient—naïve use often underperforms. Our primitives apply structured, tree-pattern, and collective scheduling to fully exploit DSMEM, enabling practical performance gains in real LLM workloads. \n7. **Clarity in Presentation of Limitations.** As mentioned in **3. About FFN.**, our on-chip fusion is \"split\" when there are large-scale data dependencies such as all-to-one reductions that cannot be fully handled within DSMEM due to capacity constraints. In these cases, our implementation falls back to global memory synchronization. However, such cases are rare in typical LLM inference workloads and the same data dependency patterns would also cause global memory traffic in baseline frameworks. Therefore, our method still maintains a performance advantage. We will make this behavior more explicit in the final version.\n8. **Depth of Algorithmic Detail.** We appreciate the reviewer’s suggestion. While our current pseudocode omits certain corner cases, our implementation can readily accommodate such scenarios. For example, load balancing can be enhanced by adaptive task-to-SM mapping in the start of the kernel, and DSMEM access patterns can be tuned in our primitive template to avoid bank conflicts. These extensions can be integrated with minimal change, thanks to the modular and general nature of our collective abstraction. As for formal performance model , We will include a roofline model that guide applicability across different workload sizes in the final version to improve clarity.\n\n[1] AMD \"RDNA4\" Instruction Set Architecture.\n\n[2] Scaling Deep Learning Computation over the Inter-core Connected Intelligence Processor with T10, SOSP 2024.\n\n[3] WaferLLM: Large Language Model Inference at Wafer Scale, OSDI 2025.\n\n[4] SGLang Blog, torchcompile-latency-optimizations, 2024-09-04.\n\n[5] vLLM Blog, torch.compile and Piecewise CUDA Graphs, 2025/01/27.\n\n[6] TensorRT-LLM Doc, Github.\n\n[7] MLC-LLM Doc, Github."}, {"review_id": "kdnAccETAA", "reviewer": "Reviewer_NWDf", "summary": "This paper proposes ClusterFusion, a new system to make large language models (LLMs) run faster during inference (when generating tokens) by fusing multiple steps—like QKV projection, attention, and output projection—into a single, optimized computation block.\n\nThe key problem is that current LLM systems run these steps separately and move data to and from off-chip memory, which slows things down. Modern GPUs like NVIDIA Hopper offer on-chip memory and fast interconnects within thread block clusters, but these are hard to use efficiently due to low-level programming.\n\nTo solve this, the authors design two new primitives:\n- ClusterReduce: for reducing values (e.g., sum or max) across thread blocks.\n- ClusterGather: for sharing data between thread blocks.\nThese allow on-chip communication and help fuse multiple operators together in a more coordinated way. The authors implement these ideas into an inference framework called ClusterFusion, which they test on Llama2-7B and DeepSeek-V2-Lite using an NVIDIA H100 GPU. Compared to popular systems like vLLM, SGLang, and TensorRT-LLM, ClusterFusion achieves an average speedup of ~1.6× in decoding time, especially on long sequences.", "questions": "Questions & Suggestions for the Authors\n\n1. How portable is ClusterFusion beyond NVIDIA Hopper?\nYour method relies on Hopper-specific features, such as DSMEM and thread block clusters. Can your approach be adapted to other GPUs, e.g., A100?\n2. What does integration with real-world LLM systems look like?\nYou claim speedups over SOTA frameworks like vLLM and TensorRT-LLM—but how would someone actually adopt ClusterFusion in production? Do you require kernel-level rewrites, custom graph compilers, or modified CUDA runtimes? \n3. Why is fusion limited to three modules (QKV, Attention, Output)?\nYou call it “ClusterFusion,” but the fusion scope still ends after Output Projection. Why not include FFN, RMSNorm, or RotaryEmbeddings? Are there hardware barriers or software trade-offs you can explain more clearly? This would clarify how general and extensible the method really is.\n4. Do you expect ClusterFusion to scale to batch sizes >1 or mixed-length prompts?\nAll your evaluations use batch size = 1. That’s common in interactive inference, but real serving systems also deal with batching and variable-length requests. Does your method still deliver speedup in those conditions? If not, could you clarify when it's useful vs. not useful?\n5. How do you recommend users tune cluster size and layout in practice?\nYou provide detailed analysis of cluster size trade-offs, but no guidance on how a practitioner should make these decisions automatically. Could you expose this as a tunable parameter? Or integrate it into a profiling-based auto-tuner?", "limitations": "1. Highly hardware-specific design\nThe entire framework is tightly coupled to NVIDIA Hopper GPUs. There is no discussion about compatibility with previous-generation GPUs (e.g., A100) or non-NVIDIA hardware (e.g., AMD, Intel).\nThe method risks becoming obsolete or irrelevant outside the Hopper ecosystem. Suggestion: Clarify how general the design is, and whether a fallback mode or alternative backend is possible.\n2. No discussion on compiler or framework support\nAlthough the paper presents an elegant low-level optimization, there is no clear path to integrating it into real-world ML compilers (like TensorRT, TVM, or PyTorch). Discuss whether a practical implementation could be exposed via kernel templates or integration with tools like Triton or CUTLASS.\n3. Limited task and batch settings\nAll experiments are conducted with batch size = 1 and on specific LLMs (Llama2-7B, DeepSeek-V2-Lite). It’s unclear whether the speedups hold in multi-batch, mixed-length, or multitask environments, which are common in actual deployments. Suggestion: Discuss potential performance changes under more realistic workloads.\n4. Potential for energy-performance tradeoffs is not explored\nWhile the work optimizes latency, it does not address power or energy use, which is critical in datacenter-scale inference. Suggestion: Even if not measured, a brief comment on energy efficiency (e.g., due to reduced off-chip memory traffic) would strengthen the story.", "rating": 4, "confidence": 4, "soundness": null, "presentation": null, "contribution": null, "strengths": "", "weaknesses": "", "quality": 3, "clarity": 3, "significance": 3, "originality": 3, "strengths_and_weaknesses": "Strengths \n\n1. Originality: Proposing structured, cluster-level primitives for intra-block GPU communication is non-trivial and novel; Extending fusion beyond single kernels via coordinated scheduling of communication and computation is a clear step forward for system-level optimization in LLM serving.\n2. Significance: This is not just a toy speedup: the ~1.6× end-to-end decoding speedup is meaningful and relevant to real-world LLM deployment at scale; The ideas are hardware-aware and reflect a strong grasp of actual GPU architecture, making them useful beyond just academic interest.\n3. Clarity: The paper is mostly clear and readable. Diagrams (especially Fig. 6 and 7) help explain the concepts well; The structure (problem → primitives → fused dataflow → results) is logical and easy to follow.\n\nWeaknesses \n\n1. Over-reliance on NVIDIA Hopper (H100):\nThe entire framework depends heavily on low-level Hopper-only features (e.g., DSMEM, thread block clusters).\n The generalization to other GPUs (e.g., A100, non-NVIDIA) is not discussed at all. This makes the solution less broadly applicable, and the paper may quickly become obsolete as hardware evolves. \n2. Fusion scope is still bounded by hardware (16-thread-block limit):\nWhile the paper claims to \"expand\" fusion, it’s still limited by the current hardware cluster size.\nThere is no strong roadmap for scaling this beyond what Hopper already allows. \n3. The method is system-level but lacks co-design with compiler/runtime stack:\nThere is no integration or discussion of compiler-level scheduling, which is critical for making this approach portable or usable at scale. How would a user actually integrate this into their LLM stack? Would HuggingFace or TensorRT need to be rewritten?\n4. Ablation and analysis are shallow in parts:\nThe ablation on DSMEM (e.g., Fig. 13) is helpful but too limited. What happens to performance if you scale to batch sizes >1, or mix prompt lengths? A stronger analysis of latency breakdowns, kernel utilization, or memory pressure would help explain where the 1.6× speedup comes from in more detail.", "was_revised": false, "final_justification": "", "rebuttal": "Thank you for your thoughtful feedback and suggestions on our paper. We would like to provide more clarification.\n\nTo answer your concerns:\n1. **Generality on other hardwares.** The techniques and rationale implemented in ClusterFusion are not limited to NVIDIA Hopper GPUs. Our fundamental goal is to enable holistic dataflow design at the larger scale. As discussed in the background section, existing fragmented design makes it difficult to achieve globally optimal. We construct ClusterReduce and ClusterGather that enable on-chip cooperation between thread blocks. On A100 GPUs, where DSMEM is not available, our on-chip collective primitives are not supported. However, as shown in Section 4.3 and Figure 13, we implement ClusterReduce and ClusterGather using global memory as a fallback, and still achieve competitive performance due to reduced kernel launches and improved holistic dataflow design. Below are the TPOT (ms) of Llama2-7B on A100:\n\n| Seqlen\\Framework | SGLang | vLLM | ClusterFusion |\n|------|------|------|------|\n| 4k | 16.8   | 16.5 | **14.1** |\n| 8k | 17.1   | 16.9 | **14.9** |\n\nMoreover, DSMEM-like inter-core communication is becoming a mainstream direction in modern architectures[1,2,3,4]. These trends indicate that ClusterFusion’s primitives and dataflow strategies are compatible with emerging hardware. Also, the primitives and holistic dataflow can be mapped to other platforms. On AMD GPUs[1], SMs, shared memory, and DSMEM used in our algorithm can be naturally mapped to compute units, local data share (LDS), and LDS-to-LDS interconnects (e.g., in RDNA 4). On dataflow architectures such as Graphcore IPUs[3] and Cerebras Wafer-Scale Engines[4], our fine-grained task partitioning strategy remains applicable across the many-core compute fabric of these systems. The holistic dataflow design philosophy as well as our inter-core tree-pattern ClusterReduce and ClusterGather can be further explored on these platforms to fully exploit their on-chip interconnects and distributed memory hierarchy. We are actively collaborating with engineers to implement our primitives on other GPU like MI300X/MI308X GPUs, enabling the same holistic, end-to-end dataflow optimization.\n\n2. **Integration into frameworks/compilers.** We have already integrated our fused kernels into the PyTorch stack via pybind. This is foundation for future integration into real production frameworks such as SGLang, vLLM, and TensorRT-LLM. For the evaluation shown in our paper, we have already use our wrapped PyTorch-compatible interfaces. These interfaces are aligned with existing kernel stack and support drop-in replacement. For framework integration, we need to replace the original kernels used in other frameworks with our ClusterFusion kernel but no CUDA runtime modifications. For CUTLASS integration, we can append our templates of ClusterReduce/ClusterGather primitives into CUTLASS template library. For ML compilers integration, we need to add a cluster pass which automatically group some thread blocks into a thread block cluster according to the data dependency for reduce/gather.\n3. **About FFN, RMSNorm and RoPE.** We have extended our primitives to fuse other modules like FFN, and the corresponding implementation(codes) is included in our supplementary material. In short, FFN fusion is feasible, but the performance gain is limited. \nFirstly, the FFN kernel must be separated from the Output projection, because the output projection in Attention introduces a all-to-one data dependency: all thread blocks must reduce their output to the same address. This is not feasible to resolve fully on-chip via DSMEM without falling back to global memory, as the limited cluster size. Secondly, we designed a fused FFN kernel where we partition the intermediate size across blocks within a cluster, and split the output dimension across different clusters. ClusterReduce is then used to compute partial sums across blocks on chip, followed by local element-wise SiLU and residual addition. However, due to the large intermediate dimensions in most mainstream LLMs (e.g., 11008 or higher), each thread block must handle a huge reduction payload. As shown in our profiling results (Section 2.3 and Figure 5), the limited DSMEM bandwidth in this case prevents us from gaining significant on-chip communication benefits. However, FFN fusion still helps reduce kernel launch overhead and global memory traffic and achieves comparable latency to existing frameworks:\n\n| Framework         | FFN Latency (us) |\n|-------------------|------------------|\n| SGLang            | 115              |\n| vLLM              | 113              |\n| TensorRT-LLM      | 107              |\n| MLC-LLM           | 110              |\n| Ours  | **118**        |\n\nRegarding RMSNorm and RoPE, these components are omitted from the paper for clarity, but are indeed included in our fused kernels. Since they can be naturally fused using shared memory or registers without requiring ClusterReduce or ClusterGather primitives. For simplicity, we did not illustrate them in the figures, but their full implementations are provided in the supplementary material.\n\n4. **The performance of multi-batch and variable length.** ClusterFusion delivers speedup across both batch size = 1 and batch size > 1 scenarios. For batch size > 1, we provide empirical results in the Appendix of the supplementary material, where we observe an average end-to-end speedup of 1.2×. As for variable-length prompts, supporting fully dynamic batching requires extending the attention module to handle runtime index computation and layout alignment. This extension is technically straightforward and will not compromise the performance benefits of our fused design, since the core communication and computation patterns remain unchanged. We are actively working on this enhancement extending our implementation to better support these use cases.\n5. **Tuner usage.** We expose cluster_size as a tunable parameter for users and offline auto-tuning. In our current implementation, cluster_size is a template parameter at the kernel level. Based on our profiling (Section 2.3), larger values like 16 are generally discouraged. So, we provide a simple tuner that performs offline benchmarking over a small candidate set (e.g., 2, 4, 8). The tuner selects the best-performing configuration for a given input models.\n6. **Energy optimizations.** Thanks for your suggention about discussion of energy-performance tradeoff. In section 4.2 and Figure 12, we report the reduced off-chip memory traffic dut to ClusterFusion optimizaions. These changes naturally reduce the frequency of global memory accesses and kernel launches, both of which are known to contribute to energy consumption on modern GPUs. We will include a discussion of these implications in the final version.\n\n[1] AMD \"RDNA4\" Instruction Set Architecture.\n\n[2] NVIDIA Blackwell Architecture Technical Brief.\n\n[2] Scaling Deep Learning Computation over the Inter-core Connected Intelligence Processor with T10, SOSP 2024.\n\n[3] WaferLLM: Large Language Model Inference at Wafer Scale, OSDI 2025."}, {"review_id": "3Iz7yLz1O1", "reviewer": "Reviewer_KUk8", "summary": "On NVIDIA Hopper GPUs, distributed shared memory is introduced to allow two thread blocks within the same thread block cluster to access each other’s shared memory. This feature provides a more efficient method for inter-block communication compared to using global memory. This paper proposes two cluster-level collective primitives that leverage this low-level feature. Based on the proposed primitives, the authors also implement a fused attention mechanism (including QKV projections, self-attention, and output projection) in a single kernel. Experiments are conducted to demonstrate the effectiveness of the proposed primitives and the fused kernel.", "questions": "Thank you to the authors for submitting to NeurIPS 2025. The paper is well written and easy to follow. Overall, I am positive about the submission. I have two concerns that I hope can help improve the paper.\n\n1. **The proposed primitives themselves do not present much novelty. Reduce and gather are standard collective primitives.**\n   While reduce and gather are standard primitives, building them using low-level PTX instructions may be a novel aspect of this work. Therefore, I recommend clarifying which low-level instructions are used (e.g., `ld.shared::cluster.u32`), possibly in a table, and explaining why they are difficult to use directly.\n\n2. **More evidence or experiments are needed to justify the necessity of the proposed primitives for the fusion.**\n   Figure 7 illustrates the fused attention operation using the proposed collective primitives. However, I did not see a fundamental limitation preventing similar kernel without using cluster-level primitives. For example, could the same thread block handle the tasks of both TB0 and TB1 in Figure 7 by reducing the tile size? If so, additional experiments comparing the block cluster approach with a single thread block would be helpful. If not, it would be beneficial to explain why, e.g., due to insufficient shared memory, with concrete numbers.", "limitations": "No", "rating": 5, "confidence": 4, "soundness": null, "presentation": null, "contribution": null, "strengths": "", "weaknesses": "", "quality": 3, "clarity": 3, "significance": 2, "originality": 3, "strengths_and_weaknesses": "**Strengths**\n\n1. The proposed techniques target important workloads (i.e., efficient LLM inference) on relevant hardware (i.e., recent NVIDIA GPUs).\n2. The results are promising.\n3. The paper is well written and easy to follow.\n\n**Weaknesses**\n\n1. The proposed primitives themselves do not present much novelty. Reduce and gather are standard collective primitives.\n2. More evidence or experiments are needed to justify the necessity of using the proposed primitives for the fusion.", "was_revised": true, "final_justification": "Thanks the authors for the detailed response which addressed majority of my concerns, increasing the score from 4 to 5.", "rebuttal": "Thank you for your thoughtful feedback and suggestions on our paper. We would like to provide more clarification.\n\nTo answer your concerns:\n1. **About novelty and the low-level PTX.** While reduce and gather are common primitives in distributed system, they are rarely considered in chip-scale with interconnect. This enable new optimization scope and opportunities, and our novelty mainly lies in the exploration of on-chip dataflow via interconnect. On Hopper GPUs, blocks within a cluster are fully connected via DSMEM, as illustrated in Figure 4 and Figure 6. And as shown in Section 2.3 and Figure 5, DSMEM performance is non-trivial and not automatically efficient—naïve use often underperforms. Our primitives apply structured, tree-pattern, and collective scheduling to fully exploit DSMEM, enabling practical performance gains in real LLM workloads. And we have provided our codes in the supplementary material. We use the following set of PTX instructions in our implementation of ClusterReduce and ClusterGather:\n\n| High-level Function | PTX Instruction | Description |\n|---------------|---------------|---------------|\n| Init DSMEM barrier                  | `mbarrier.init.shared::cta.b64`                                | Initialize inter-block synchronization barrier within a CTA                |\n| Notify barrier with transfer data size    | `mbarrier.arrive.expect_tx.shared::cta.b64`                    | Notify barrier with the expected data size of Reduce/Gather           |\n| Map shared memory across blocks     | `mapa.shared::cluster.u32`                                     | Map the shared memory address from one thread block to another             |\n| DSMEM copy | `cp.async.bulk.shared::cluster.shared::cta.mbarrier::complete_tx::bytes` | Copy data across shared memory of blocks with barrier synchronization      |\n| Wait for DSMEM transfer completion | `mbarrier.try_wait.parity.shared::cta.b64`                     | Used by the consumer thread block to detect whether the producer has completed the DSMEM write            |\n\n\nAll of these instructions are necessary for implementing our ClusterReduce and ClusterGather primitives. While these instructions are documented in PTX ISA references [1,2], they are low-level, semantically fragmented, and designed only for fine-grained behaviors (e.g., address mapping or signaling). Their correct composition into usable patterns is non-trivial. Many of them take complicated parameters (e.g., barrier state tokens, data sizes, memory space specifiers) and require strict ordering to avoid deadlocks or incorrect synchronization. Our contribution is to encapsulate these instructions into robust, reusable cluster-level primitives. This abstraction significantly reduces development cost and makes it practical to integrate DSMEM communication into LLM inference frameworks.\n\n2. **The necessity of the proposed primitives for fusion.** Thanks for your insightful suggestion. Our cluster-level primitives are necessary to scale fusion across realistic LLM workloads with high parallelism. Indeed, it is possible to let a single thread block handle the tasks of both TB0 and TB1 in Figure 7 (equivalent to setting cluster size = 1), which eliminates the need for inter-block communication. We have implemented and evaluated this variant. The evaluation result of ClusterFusion core kernel latency(us) under different cluster size are:\n\n| Seqlen\\Cluster Size | 4 | 2 | 1 (w/o Cluster-level Primitives) |\n|------|------|------|------|\n| 4k | 108 | 121 | **212** |\n| 8k | 125   | 132  | **224** |\n\nAs above results show, using only one thread block significantly degrades performance. This is because a single block often lacks sufficient parallelism. In practice, we need to split the workload across more thread blocks (e.g., along the head or token dimension to fully utilize the hardware). However, such splits introduce inter-block data dependencies that require reduction or aggregation. This is where our on-chip ClusterReduce and ClusterGather primitives become critical: they enable efficient intra-cluster communication without falling back to global memory. \n\n\n[1] Parallel Thread Execution ISA Version 8.8, NVIDIA.\n\n[2] NVIDIA H100 Tensor Core GPU Architecture."}, {"review_id": "6TBtyp7ucu", "reviewer": "Reviewer_1tSz", "summary": "This paper applied the idea of \"collective communication\" or tensor parallellism/sharding between multiple devices to the sharding/on-chip collectives between threading blocks/Streaming Processors within a single device, and demonstrates it enables kernel fusion of the entire self-Attention module and performance gain over SOTA inference kernels/engines.", "questions": "1. It is a bit confusing to use the term \"Cluster\" as it typically refers to distributed multi-device/node system. While the core ideas of collectives are the same with a cluster of devices and a cluster of threadblocks, it may be a bit misleading for readers. Similarly the authors called the two primitives Cluster-Reduce/Gather, aren't they just Block/SM level all-reduce/gather?\n\n2. The collectives are implemented via tree pattern, have you considered ring pattern?\n\n3. How do you define the bandwidth between a cluster of SMs? is it defined by peer-to-peer bw if so I am not sure why it is affected by cluster size or is it deifned by bipartite or all-reduce bw etc?\n\n4. Why would the active SMs descrease with size of clusters?\n\n5. Line 256 mentioned the kernel launch overhead is reduced by an order, however based on Figure 7, the kernel is a fusion of 4 kernels of Figure 3, how is it able to reduce kernel launch overhead by an order (or 10 times)?", "limitations": "* While included in appendix, I think the authors should discuss the efficacy of the method on higher batchsize > 1 in the main body. e.g., why does higher batch size see a much less performance gain over baselines when no weight is involved (unlike MLPs) so off-chip traffic is purely proportional to batchsize\n\n* The onchip collectives rely on creating buffers with onchip shared-memory, what if the model/kernel is too large so it runs out of the pretious onchip memory?", "rating": 5, "confidence": 5, "soundness": null, "presentation": null, "contribution": null, "strengths": "", "weaknesses": "", "quality": 3, "clarity": 3, "significance": 3, "originality": 3, "strengths_and_weaknesses": "Strength:\n* The motivaiton is intuitive and novel on GPU programming: modern LLM decoding on GPU is bottlenecked by off-chip memory bandwidth and kernel launch/synchronization overhead. And the onchip collective communication between SMs/Thread-blocks and reduce off-chip traffic and promote kernel fusions\n\n* The author abstract and implemented two effective collective communications among a thread block cluster: All-Reduce and All-Gather, which serves the fundamental to perform tensor/operator sharding onchip\n\n* The authors demonstrated significant performance gain e.g., 50% over SOTA inference engine/kernels with BS=1\n\nWeakness:\n* While relative new among GPU, onchip datalfow/collectives have been widely explored by dataflow architecture accelerators like cited GraphCore/Cerebras etc, the associated aggressive kernel fusion has also been studied by https://arxiv.org/pdf/2410.23668 which extends kernel fusion to the entire transformer model. Recently there is also research about mega kernel fusion on GPU: https://hazyresearch.stanford.edu/blog/2025-05-27-no-bubbles, which should be included in discussion\n\n- The authors didn't discuss enough limitations of this approach, detailed below", "was_revised": true, "final_justification": "The rebuttal addressed most of my questions with detailed explanations.", "rebuttal": "Thank you for your thoughtful feedback and suggestions on our paper. We would like to provide more clarification.\n\nTo answer your concerns:\n1. **About the term \"Cluster\".** Sorry about the confusion, we will use the term \"SM Cluster\" to avoid this confusion. Indeed, our primitives operate at the block/SM level. We chose the term “Cluster” to align with NVIDIA’s terminology on Hopper GPUs[1,2], where a cluster refers to a group of thread blocks with shared memory connectivity. \n2. **Different communication pattern.** Thanks for your point. We have carefully considered both the tree pattern and the ring pattern (ScatterReduce + AllGather), and we chose the tree pattern because it requires fewer communication rounds, which leads to lower latency when the data size is relatively small. \nWe compared the two patterns both theoretically and empirically under cluster size = $N$ and data size = $K$. In theory, the ring pattern requires $(N-1)$ rounds for ScatterReduce and another $(N-1)$ rounds for AllGather, totaling $2(N - 1)$ communication rounds, with each round transferring $\\frac{K}{N}$ data per node. In contrast, the tree pattern only needs $\\log_2 N$ rounds, with each round transferring $K$ data per node, as described in our paper. \nThe following table summarizes the theoretical communication characteristics for both patterns:\n\n| Pattern     | Comm Rounds     | Data per Round | Total Data per Node |\n|-------------|------------------|----------------|----------------------|\n| Ring        | 2(N - 1)         | K / N          | 2(N - 1) * (K / N)   |\n| Tree        | log₂(N)          | K              | K * log₂(N)          |\n\nFor example, when $N=2$ and $N=4$:\n\n| Pattern \\ N | Rounds | Data/Round | Total Data |\n|-------------|--------|------------|-------------|\n| Ring (N=2)  | 2      | K/2        | K           |\n| Tree (N=2)  | 1      | K          | K           |\n| Ring (N=4)  | 6      | K/4        | 1.5K        |\n| Tree (N=4)  | 2      | K          | 2K          |\n\nAlthough the tree pattern transfers more total data, its fewer communication rounds are more favorable for latency because the data size $K$ is small (64–128 bytes) in the ClusterFusion dataflow and DSMEM bandwidth is not fully utilized. This is also reflected in our empirical results, where we report the kernel latency(us) of the core ClusterFusion kernel:\n| Pattern\\Cluster Size | 2 | 4 |\n|------|------|------|\n| Ring | 135 | 127 |\n| Tree | 132   | 125 |\n\nWe plan to include a more detailed discussion about different communication pattern in the final version. \n3. **Definition of the bandwidth between SMs within cluster.** The bandwidth shown in Section 2.3 and Figure 5 refers to peer-to-peer bandwidth. As cluster size increases, peer-to-peer bandwidth decreases. Based on the patent and paper referenced in that section, the interconnect topology for DSMEM in Hopper GPUs is likely a crossbar or crossbar-like architecture, which allows any SM to communicate with any other within a cluster. While a crossbar provides full connectivity, it does not provide dedicated bandwidth for all possible communication pairs simultaneously. Instead, the internal paths of the crossbar must be shared among all active transfers, and contention arises as more thread blocks in a cluster attempt concurrent communication. \n\n4. **The reason for the decreased active SMs.** To the best knowledge of us based on public architectural information[1,2,3], the number of active SMs may decrease with larger cluster sizes due to hardware-level SM allocation constraints. Whenever the cluster size exceeds 2, the hardware grouping may lead to SM underutilization due to yield-related SM inactivation and cluster alignment issues. On H100 SXM GPUs, the 132 SMs are not evenly distributed: three GPCs[2] have 18 SMs, four have 16, and the last one contains 8 + 6 SMs, where the last 6 are special. For example, when using Cluster Size = 4, this grouping cannot fully utilize all SMs: in the first three GPCs, 2 SMs per GPC are left out (3 × 2 = 6), and in the last GPC, all 6 special SMs are likely excluded, totaling 12 unused SMs. This leaves only 120 active SMs which aligns with our observed drop in active SMs during profiling.\n5. **About kernel launch overhead in Figure 7.** Thanks for pointing this out. Actually, Figure 7 shows the total kernel launch overhead of the end-to-end decoding. Our fused kernel reduces the overall kernel launch overhead across all layers. And our kernel-level launch overhead is indeed one order of magnitude lower than the CUDA Graph launch overhead used in existing frameworks. For example, in the Llama2-7B model, the end-to-end kernel launch overhead drops from 1.057 ms in the baseline to 0.169 ms with ClusterFusion.\n6. **The performance of multi-batch.** For batch sizes >1, we observe an average 1.2× end-to-end speedup in the Appendix provided in the supplementary material. As batch size grows, the workload becomes more compute-intensive and less memory-bound, reducing the optimization space for memory traffic. Although the off-chip memory traffic increases in higher batch size, the proportion of activation memory traffic and sync decreases due to the significantly increased compute intensity. Nevertheless, ClusterFusion still achieves meaningful gains, demonstrating its robustness across diverse inference scenarios. \n7. **The reliance on shared memory buffer.** In our current implementation, the algorithm relies on either reusing existing shared memory tiles (e.g., weight tiles) or allocating a small additional buffer. This is necessary because our reduction is single-phase: the data being transmitted and the data being reduced in each iteration reside in the same memory region, requiring intermediate buffering to avoid conflicts. However, when shared memory becomes insufficient, we can switch to buffer-free variants such as ScatterReduce. In this approach, each thread block writes its data to non-overlapping destinations, allowing communication and reduction to proceed simultaneously without intermediate buffering. \n8. **Comparison with other concurrent work.** Thank you for pointing out two recent, concurrent work: Kernel Looping[4] and MegaKernel[5](publicly available after our submission). Kernel Looping is designed for a dataflow architecture accelerator—specifically the SambaNova SN40L Reconfigurable Dataflow Unit (RDU). It leverages the repetitive nature of decoder layers in LLM to fuse different layers, reducing multiple kernel launches to a single launch. This approach serves a role similar to CUDA Graphs in the GPU. Our method also significantly reduces kernel launch overhead on GPUs, and we empirically outperform frameworks that already leverage CUDA Graphs. The motivation of MegaKernel is similar to ours in that both aim to reduce memory bubble. Their approach is orthogonal to our work. They focus on fine-grained shared memory reuse and overlapping activation storage with weight loading. Despite this, MegaKernel still relies on global memory for intermediate result synchronization, whereas our work introduces cluster-level primitives to enable fast, on-chip synchronization. We will incorporate a detailed discussion of these recent works in the related work of the final paper.\n\n[1] NVIDIA CUDA C++ Programming Guide.\n\n[2] NVIDIA H100 Tensor Core GPU Architecture.\n\n[3] Benchmarking Thread Block Cluster, HPEC 2024.\n\n[4] Kernel Looping: Eliminating Synchronization Boundaries for Peak Inference Performance, arxiv.\n\n[5] Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B, Hazy Research."}]

Author rebuttal:

Meta-review: This paper proposes ClusterFusion, a new system that accelerates LLM inference by fusing multiple steps—such as QKV projection, attention, and output projection—into a single optimized computation block. All reviewers agree that this work makes solid contributions, and the AC recommends acceptance.
Decision: Accept (poster)