Fused Generation Kernels: Methods & Impacts

• Fused generation kernels are compute kernels created by automatically fusing multiple operations (e.g., map, reduce, GEMM) to minimize memory traffic and kernel launch overhead.
• They employ techniques such as dependence analysis, fusion plan enumeration, and resource-constrained scheduling to optimize data reuse across registers and on-chip buffers.
• Empirical results show significant speedups (up to 4× or more) in domains like deep learning and scientific simulations, although trade-offs in occupancy and code size must be managed.

Fused generation kernels are GPU or CPU compute kernels produced by the automatic or semi-automatic fusion of multiple computational stages—typically map, reduce, matrix-matrix/vector-matrix, or more complex dataflow graph operations—into a single, monolithic kernel launch. This fusion minimizes global-memory traffic by passing intermediate data via registers or on-chip buffers, reduces kernel launch overhead, and enables co-optimized scheduling or local synchronization. Fused generation is crucial for achieving high arithmetic intensity and exploiting on-chip resource hierarchies in bandwidth-bound settings such as deep learning, scientific simulation, graph analytics, and distributed generation tasks (Filipovič et al., 2013, Long et al., 2018, Huang et al., 15 Dec 2025, Zhang et al., 12 Feb 2026, Hoque et al., 2024).

1. Fundamental Principles and Motivations

Kernel fusion originated as a direct response to the memory–compute gap in modern accelerators: floating-point throughput increases have outpaced improvements in DRAM bandwidth, exposing memory bandwidth as the primary bottleneck for chaining multiple compute stages. By fusing kernels, intermediate values are retained in registers or shared memory rather than written out and reloaded from global memory, realizing up to 2–4× reduction in memory traffic for common patterns (e.g., map-map, map-reduce, GEMM+pointwise) (Filipovič et al., 2013, Adnan et al., 2015, Zhang et al., 12 Feb 2026).

There are two overarching forms:

• Vertical Fusion (VF): sequential stages (producer-consumer) are merged so intermediates never leave chip.
• Horizontal Fusion (HF): independent launches over distinct input slices are merged by extending the grid (e.g., blockIdx.z over batch) (Li et al., 2020, Amoros et al., 9 Aug 2025).

Both forms aim to maximize arithmetic intensity $AI=F/G$, where $F$ is the number of floating-point operations and $G$ is the bytes of global memory traffic, favoring compute-bound execution for bandwidth-limited operators (Filipovič et al., 2013, Zhang et al., 12 Feb 2026, Zuo et al., 22 Apr 2026).

2. Algorithmic and Compiler Techniques for Fusion

Fused generation kernel construction is typically automated by a source-to-source compiler, code generator, or metaprogramming system:

• Dependence Analysis: The input is either a sequence of elementary library calls (e.g., BLAS routines), a computation DAG (as seen in XLA or MLIR), or a high-level API (e.g., Fused Kernel Library). The compiler builds a data-dependency DAG where nodes represent compute stages (“elementary functions” or ops) and edges represent data flow or synchronization points (Filipovič et al., 2013, Long et al., 2018, Amoros et al., 9 Aug 2025).
• Fusion Plan Enumeration: The system enumerates all maximal fusible subgraphs without breaking required global barriers, atomic steps, or on-chip mapping constraints. For map-reduce or block-wise reductions, fusion is kept within operable block or thread boundaries to avoid introducing illegal synchronizations (Filipovič et al., 2013, Long et al., 2018).
• Resource-Constrained Tiling and Scheduling: An empirical or analytic cost model predicts kernel performance parameters given tile sizes, shared/register usage, and occupancy. Modern frameworks like FlashFuser integrate distributed shared memory (DSM) as an additional level in the memory hierarchy and extend the search space to utilize intra-cluster high-bandwidth memory with cost models quantifying data movement at each level (Huang et al., 15 Dec 2025).
• Code Generation and Inlining: The selected fusion schedule is emitted as a single CUDA/C++ kernel (or similar), inlining load, compute, store, and synchronization calls. If needed, partial barriers (block-level or sub-block via PTX bar.sync) are inserted precisely where the thread-to-data mapping changes (Li et al., 2020, Amoros et al., 9 Aug 2025).
• Cross-Framework and Distributed Fusion: Systems like Diffuse raise task and kernel fusion to the distributed level, leveraging a scale-free IR and JIT-compilation based on partition privilege and launch domain dependency to produce fused kernels across multi-GPU systems (Yadav et al., 2024).

3. Concrete Transformation Patterns and Examples

BLAS-Style Map/Reduce Fusion

Typical map-map and map-reduce transformations replace two sequential launches and a global memory round-trip with a single kernel launch. E.g.,

• Map-Map:

$$A' = \mathrm{map}(f, X),\quad B' = \mathrm{map}(g, A')$$

fused as

$C = \mathrm{map}(g\circ f, X)$

with data passed via registers or shared RAM, guarded by thread/block mapping compatibility (Filipovič et al., 2013).

• Map-Reduce: Per-block reductions fuse map and reduction such that partial block-level results remain on-chip, while the global reduction is completed by a separate kernel or via atomics (Filipovič et al., 2013).

Fused Quantized MatMul

In transformer inference, dequantization and GEMM are fused to avoid materializing full-precision weights. For W4A16, packed 4-bit weights are dequantized in registers and directly accumulated, with SplitK decomposition used for parallelization (Hoque et al., 2024).

Deep Fusion of Transformer Subgraphs

DeepFusionKernel implements entire SwiGLU blocks (GEMM-Up, GEMM-Gate, SiLU, gating, GEMM-Down) as a single kernel. Tiles of activations and weights are loaded exactly once into registers/shared memory, and interleaved GEMM and pointwise compute chains maximize cache reuse and minimize HBM traffic, yielding up to 13.2% speedup on H100 (Zhang et al., 12 Feb 2026).

Kernel Fusion of Sparse and Graph Operations

FusedMM unifies sparse sampled-dense-dense matrix multiplication (SDDMM) and sparse-dense (SpMM) using a programmable five-stage abstraction (VOP, ROP, SOP, MOP, AOP) and generates SIMD- and memory-efficient fused kernels (Rahman et al., 2020).

Fused Generation in Scientific/HPC and Video Analytics

In hyperbolic diffusion and video feature tracking, fusion cuts multiple full-array passes down to one, replaces global intermediates by on-chip circular buffers, and achieves consistent 2–4× speedups (Trojak et al., 2021, Sewall et al., 2017, Adnan et al., 2015).

4. Hardware and Memory Hierarchy Considerations

Recent architectures have deep on-chip hierarchies:

• Registers: fastest, but limited in size. Used for immediate intermediates.
• Shared Memory (SMEM, L1): on-SM, used for intra-block staging.
• Distributed Shared Memory (DSM, L1.5 on Hopper/H100): connects SMs within a cluster, supports inter-block on-chip collectives (reduce, shuffle, scatter) (Huang et al., 15 Dec 2025).
• L2 Cache and HBM/DRAM: accessed only for kernel input/output or when on-chip space is exceeded.

FlashFuser and similar tools model data movement explicitly across all levels: $C_l = V_l / B_l$ where $V_l$ is the data volume and $B_l$ is the bandwidth at level $l$. The optimal fusion plan minimizes $\max_l C_l$, subject to per-level capacity constraints (Huang et al., 15 Dec 2025).

Fused generation kernels thus aim to maximize reuse of data in registers/SMEM/DSM, only spilling to L2/DRAM when unavoidable.

5. Performance Impact and Empirical Results

Empirical studies demonstrate substantial speedups, with kernel fusion moving many workloads from bandwidth-bound to compute-bound regimes. Selected numeric results include:

Domain	Reported Kernel Speedup	End-to-End Speedup	Reference
BLAS-1/2 (AXPYDOT, etc.)	up to 2.61×	n/a	(Filipovič et al., 2013)
Deep learning (TF, XLA)	1.15–3.5×	up to 1.20×	(Long et al., 2018)
Triton W4A16 quantized	1.24× avg (H100 PCIe)	n/a	(Hoque et al., 2024)
SwiGLU MLP fusion	3–13%	n/a	(Zhang et al., 12 Feb 2026)
FlashFuser (DSM fusion)	up to 4.1×	1.24×	(Huang et al., 15 Dec 2025)
FairyFuse (CPU ternary)	29.6×	1.24× over Q4_K_M	(Zuo et al., 22 Apr 2026)
FKL (GPU template fusion)	up to >1000×	n/a	(Amoros et al., 9 Aug 2025)
Video tracking (K20)	2–3×	up to 40,000 FPS	(Adnan et al., 2015)
Graph embedding (FusedMM)	up to 28×	n/a	(Rahman et al., 2020)
Distributed fusion (Diffuse)	up to 10.7× (128 GPUs)	up to 1.86×	(Yadav et al., 2024)

Speedup magnitude depends on the memory–compute ratio, primitive types fused, kernel occupancy, and empirical bottlenecks such as kernel launch or interconnect overhead. Fused kernels typically achieve bandwidth utilization of 65–90% of hardware peak (GTX 480, K20, Skylake, H100), with DNN/BLAS workloads showing the strongest gains in memory-dominated settings.

6. Fusion Limitations, Trade-Offs, and Practical Concerns

• Occupancy vs. Register/Shared Pressure: Deeply fused kernels increase register and shared memory usage, potentially reducing occupancy. Tiling/autotuning mitigates this, but practical occupancy should remain above 50% for best utilization (Zhang et al., 12 Feb 2026, Hoque et al., 2024).
• Global Synchronization Barriers: Vertically fusing kernels with required grid-level barriers (e.g., final reductions, non-associative ops) is precluded by correctness constraints (Filipovič et al., 2013, Li et al., 2020).
• Code Size and Complexity: Excessive fusion can result in large monolithic kernels, code-size blowup, and lower register allocation flexibility. Most practical systems limit fusion width and depth (Li et al., 2020, Amoros et al., 9 Aug 2025).
• Pattern Coverage: Fusion techniques are effective primarily on elementwise, reduction, and regular sparse patterns. Arbitrary stencils, indirection, or side-effect-heavy code may fall outside current fusion frameworks (Filipovič et al., 2013, Sewall et al., 2017, Rahman et al., 2020).
• ISA and Portability: Some approaches (e.g., FairyFuse's ternary AVX-512 kernels) are ISA-specific and may not generalize without hardware or software emulation (Zuo et al., 22 Apr 2026).

7. Evolving Directions and Generalization

Recent advances have extended the fusion paradigm:

• Distributed and Communication-Collective Fusion: Combining local compute and remote collective operations (e.g., embedding + All-to-All, GEMV + AllReduce), with persistent kernels issuing GPU-initiated communication for fine-grained overlap (Punniyamurthy et al., 2023).
• Cross-Library/Framework Fusion: Intermediate representations (MLIR, XLA HLO) act as fusion substrates spanning entire frameworks or pipelines, supporting composable optimization across otherwise separate code boundaries (Long et al., 2018, Yadav et al., 2024).
• Meta/Compile-Time Fusion: Template-based libraries (e.g., FKL) enable per-instance compile-time fusion for multi-operation pipelines, leveraging modern C++ metaprogramming features without custom code generation (Amoros et al., 9 Aug 2025).
• Hardware-Aware Search and Autotuning: DSM-aware search (FlashFuser) integrates multiple on-chip memory levels into the fusion decision, allowing scaling beyond SMEM capacity (Huang et al., 15 Dec 2025).
• Domain-Specific Extensions: Scientific codes (HFAV for stencils, PyFR for DG/FVM schemes), graph analytics (FusedMM), and distributed dataflow pipelines all benefit from tailored fused generation compilers (Sewall et al., 2017, Trojak et al., 2021, Rahman et al., 2020, Yadav et al., 2024).

Fused generation kernels are now established across a broad set of domains as the primary vehicle for bridging the emerging bandwidth–compute gap, with rapid advances driven by hardware trends (on-chip memory, distributed fabrics) and the proliferation of high-level IRs and metaprogramming environments.