PaperlayerView exact revision on arXiv

Source-linked AI summary

SyclKittens: A Tile Programming Model for Programmers and Coding Agents on Intel GPUs

Yehong Jiang, Sheng Chen, Fangwen Fu, Yen-Kuang Chen, Xinmin Tian, Stuart H. Sul, Simran Arora

arXiv:2610.04277v1cs.DC

TL;DR

Coding agents can write and tune GPU kernels, but their performance depends on whether their interface exposes raw capabilities or encodes hardware-specific methods. SyclKittens provides reusable Intel GPU methods while keeping workload policies programmable; the paper reports faster agent-written GEMM and gains from co-designed kernels and inference.

  • Problem

    The paper asks whether a tile programming model can make Intel GPU methods reusable by programmers and coding agents while preserving workload-specific control.

    Source: §1 · p. 1

  • Method

    SyclKittens encodes Intel GPU operand layouts, cooperative cache prefetch, and topology-aware communication as reusable operations while leaving workload policies programmable.

    Source: §1 · p. 2

  • Results

    With SyclKittens, Opus reaches 82.1% of oneDNN on single-GPU GEMM, versus 57.5% with Raw SYCL; the co-designed suite approaches oneDNN on GEMM and improves inference against the stated baselines.

    Source: §1 · p. 3, §7 · p. 12

  • Takeaways & Limitations

    Engineers and agents can refine workload-specific schedules in SyclKittens while using its encoded architecture methods to build a broad kernel suite.

    Source: §1 · pp. 2–3

    Transfer to Crescent Island would require validating operand layouts and prefetch parameters and adapting communication schedules to its PCIe-based topology.

    Source: §2.2 · p. 4

Abstract

from arXiv · show

New AI accelerators arrive before the kernels that make them fast, because peak kernel performance requires architecture-specific expertise in operand pipelines and data-movement techniques. Coding agents can now write, compile, and tune kernels on their own, so they could greatly accelerate kernel development and optimization. What agents produce depends on the interfaces they are given. These interfaces may expose a machine's raw capabilities or encode the known-good methods for exploiting its hardware features efficiently. We test the performance impact of these two interfaces on Intel GPUs through controlled experiments with three coding models on three kernels. With raw SYCL and execution feedback, Opus 4.8, the strongest of the three models tested, writes a single-GPU GEMM kernel reaching only 57.5% of Intel's tuned oneDNN library. We present SyclKittens, a hardware-aware tile programming model for Intel GPUs that encodes these known-good methods, so its operations place operands in matrix-engine layouts, prefetch through the L1 cache, and communicate over the fabric that links GPU stacks into a node. With SyclKittens under the same agent, task, and feedback budget, the agent reaches 82.1%, showing that encoded methods turn hardware capabilities into performance. Workload-specific policies stay programmable, and engineers and agents jointly refine these schedules in SyclKittens to build a kernel suite that reaches ~96% of oneDNN in geometric mean across GEMM shapes. The suite runs Llama-3.1-8B inference 1.59x faster than torch$.$compile on one GPU and up to 2.91x faster than a matched multi-GPU decode path on Intel's oneCCL. SyclKittens is open source and available at https://github.com/intel/SyclKittens.

1 Introduction

The introduction asks whether an Intel GPU tile model can make hardware methods reusable by programmers and coding agents without removing workload-specific control. It presents SyclKittens and reports gains in agent-written kernels, a co-designed kernel suite, and Llama-3.1-8B inference.

1 Introduction

  • Under the same agent, task, and feedback budget, Opus raises single-GPU GEMM from 57.5% to 82.1% of oneDNN; Figure 1b shows it reaching Raw SYCL's final performance with fewer feedback calls.The controlled study isolates the effects of encoded methods, examples, and execution feedback.Source: Figure 1 · p. 2, §1 · p. 3
  • SyclKittens encodes Intel-specific operand-layout, cache-prefetch, and topology-aware communication methods as operations while leaving workload policies programmable.Programmers and agents choose tile shape, prefetch ownership and distance, and work distribution.Source: §1 · p. 2
  • Mapping register tiles from 16 × 16 to 32 × 32 raises cumulative GEMM throughput from 37.4% to 68.2% of oneDNN across 27 shapes.The design maps tiles to XMX's native execution width and fuses 2-D loads with hardware packing for operand B.Source: §1 · p. 2
  • Disjoint prefetch ownership raises tuned GEMM from 91.0% to 96.2% of oneDNN, while topology-aware gathering delivers 1.42–2.15× flat-gather distributed-GEMM throughput.The gather schedule fetches each shard once per remote package and replicates it over MDFI.Source: §1 · p. 2
  • Engineer-agent co-design yields a suite reaching ∼96% of oneDNN in geometric mean across 27 GEMM shapes and 1.21× fused SDPA with IPEX throughput on causal MHA.The suite also includes normalization, RoPE, GEMV, multi-GPU collectives, and fused distributed kernels.Source: §1 · p. 3
  • For Llama-3.1-8B inference, SyclKittens improves one-GPU full-request throughput by 1.59× over torch.compile alone and beats the oneCCL-based multi-GPU decode path by up to 2.91× on four GPUs.The one-GPU path also runs configurations where torch.compile alone runs out of memory.Source: §1 · p. 3

2 Background and Motivation

Tile programming models separate tile-level algorithms from architecture-specific layouts, movement, and scheduling. The Intel Max GPU's register, operand, and fabric characteristics shape SyclKittens, while raw-SYCL agent kernels illustrate the coordination methods the model encodes.

2.1 Tile Programming and Coding Agents

  • Tile programming models sit between tensor compilers and low-level GPU programming, separating tile algorithms from architecture-specific layouts, movement, and scheduling.ThunderKittens and HipKittens provide NVIDIA and AMD hardware-aware interfaces; SyclKittens extends the approach with Intel-specific operations and full-model inference evaluation.Source: §2.1 · p. 3

2.2 Intel GPU Model

  • The Intel Max node exposes 12 stacks as GPUs across six two-stack packages, with each stack presented as a SYCL device.The paper uses “tile” only for register or shared tiles in the language, not hardware partitions.Source: §2.2 · p. 4, Figure 2 · p. 5
  • XMX consumes 16-lane subgroups, and SyclKittens' matrix path uses L1 prefetch followed by 2-D demand loads into subgroup registers without intermediate SLM staging.These execution and transaction widths determine register-tile mapping.Source: §2.2 · p. 4
  • Topology-aware gathering crosses XeLink once per remote package and replicates over MDFI, delivering 1.42–2.15× the end-to-end distributed-GEMM throughput of flat gathering.MDFI has substantially higher bandwidth than XeLink.Source: §2.2 · p. 4

2.3 What Raw-SYCL Agents Implement

  • In single-GPU BF16 GEMM, Opus's raw-SYCL kernel reaches 57.5% of oneDNN despite using XMX, 2-D loads, VNNI layout, and prefetch.Its thread groups do not coordinate prefetching and repeatedly request the same operand tiles.Source: §2.3 · p. 4
  • Across nine correct raw-SYCL distributed-GEMM trials, no candidate combines local-output computation, XMX, and topology-aware communication; the best reaches 42.3% of the SyclKittens hand reference.The trials vary in decomposition and communication choices, including replicated-compute strategies that ignore MDFI/XeLink topology.Source: §2.3 · pp. 4–5

3 SyclKittens

SyclKittens provides Intel GPU tile operations for matrix computation, operand movement, and communication while leaving workload schedules programmable. Its layout, prefetch, and topology-aware methods support high GEMM throughput and reduce distributed-GEMM communication overhead.

3 SyclKittens

  • SyclKittens is a header-only C++ tile model whose backend lowers tile operations to SYCL or SPIR-V and selects physical lane mappings.Distributed layouts, peer memory, and GPU-resident synchronization extend the model across a node.Source: §3 · p. 5

3.1 Tile Interface

  • Register tiles represent typed two-dimensional arrays with selectable layout and physical base shape, while global layouts address tensor coordinates.Kernels choose the order of load, compute, and store stages; parallel global layouts extend the coordinate space across devices.Source: §3.1 · p. 5

3.2 Matrix Operand Layouts and Packing

  • SyclKittens maps 16-lane subgroups to XMX layouts and fuses 2-D loads with VNNI packing for the column-layout B operand.Direct L1 loading avoids the multiple ordinary reads needed to load packed B from SLM on Intel Max GPUs.Source: §3.2 · pp. 5–6

3.3 Cooperative Cache Prefetch

  • Cooperative prefetch assigns sub-tile regions to computing subgroups, letting kernels choose ownership and lookahead distance.The prefetch phase warms L1 without writing registers; a later demand load transfers cached operands into subgroup registers.Source: §3.3 · p. 6
  • 96.3% of oneDNN is reached in the cumulative GEMM ablation after layout, prefetch, and per-shape prefetch-distance changes.The six-node, 27-shape comparison reports geometric-mean throughput; the sequence begins at 37.4%.Source: §3.3 · p. 6, Table 1 · p. 6

3.4 Topology-Aware GPU Communication

  • 1.42–2.15× the end-to-end throughput of the flat gather pipeline is achieved by gathering each shard once per remote package and replicating it over MDFI.The two ratios apply to the smaller and larger shapes, respectively.Source: §3.4 · p. 6

4 Controlled Coding-Agent Evaluation

Controlled evaluations compare coding-agent kernels across programming contexts, tasks, and feedback protocols under fixed budgets. SyclKittens improves some results substantially, while effects vary by task and model; iterative refinement also improves eligible correct kernels.

4.1 Experimental Design

  • Each model-context pair receives three independent starts, and cold- and warm-iterative runs receive at most 20 evaluator calls with performance feedback.Selected candidates are graded on five unseen seeds; only correct kernels receive performance scores.Source: §4.1 · p. 7

4.2 Results

  • 82.1% of oneDNN is reached by Opus on single-GPU GEMM in cold iteration with SyclKittens, up from 57.5% with Raw.All nine Opus single-GPU trials were correct; SyclKittens + examples reached 82.6%.Source: §4.2 · p. 7, Table 3 · p. 8
  • The median Opus SyclKittens trajectory passes Raw’s final level on the second evaluator call, separately measuring search efficiency from peak@3 endpoints.Figure 1(b) reports the trajectory comparison.Source: §4.2 · p. 7
  • 79.5% for Opus with SyclKittens + examples contrasts with 18.6% for Sonnet with Raw on GQA, showing context rankings vary by model and task.Haiku’s results remain between 0.7% and 1.3% across contexts, and all 27 cold-iterative GQA trials are correct.Source: §4.2 · p. 7, Table 4 · p. 8
  • Without feedback, examples raise Opus peak@3 from 63.4% to 81.2% on single-GPU GEMM and from 36.7% to 59.4% on distributed GEMM.On distributed GEMM, correctness also rises from 2/3 to 3/3; measured example benefits depend on model, task, and starting condition.Source: §4.2 · p. 8
  • 106.6% of the distributed reference is reached by the strongest warm-iterative kernel after computing local output-row shards before gathering.All thirteen eligible distributed sources in the two SyclKittens contexts remained correct and improved, avoiding redundant full-GEMM work on every GPU.Source: §4.2 · p. 8
  • The strongest no-feedback Opus SyclKittens + examples kernel combines XMX, 2-D loads, VNNI packing, a 32×32 register base shape, and three-tile prefetch.The correct single-shot Raw kernel uses vector FMA; all contexts expose the same SPIR-V and XMX instructions.Source: §4.2 · p. 8

5 SyclKittens Inference Kernel Suite

Engineers and coding agents co-designed the SyclKittens inference suite, which evaluates GEMM, attention, and decode-critical operators against vendor paths.

5 SyclKittens Inference Kernel Suite

  • Engineers and coding agents jointly refine tile shapes, prefetch ownership and distance, and work distribution under engineer-set correctness and performance checks.Source: §5 · p. 8

5.1 Experimental Setup

  • Operator baselines dispatch through oneDNN, fused SDPA, or oneCCL; standalone decode GEMV calls oneMKL directly, and fusion and end-to-end results use torch.compile.Source: §5.1 · p. 8
  • Sweeps warm up before timing and co-measure each SyclKittens candidate with its baseline; GEMM must pass a relative Frobenius error below 0.02.Source: §5.1 · pp. 8–9

5.2 Kernel Results

  • 96.3% of oneDNN’s geometric-mean throughput across 27 GEMMs is reached by the SyclKittens GEMM, which exceeds oneDNN on three shapes and peaks at 1.17×.Source: §5.2 · p. 9
  • The GQA kernel reuses HipKittens’ tiled attention dataflow without algorithmic change, while reimplementing tile operations with Intel matrix, movement, and prefetch instructions.Source: §5.2 · p. 9
  • Fused RMSNorm, LayerNorm, and RoPE use one launch, while split-K GEMV maps output-row reductions across subgroups to keep the GPU occupied as tensor-parallel shards shrink.Source: §5.2 · p. 9

6 End-to-End Inference

The section measures SyclKittens-composed Llama-3.1-8B inference on one GPU and across a node. It reports throughput gains alongside distinctions between matched-resource comparisons and capacity results.

6 End-to-End Inference

  • 1.59× geometric-mean full-request throughput over torch.compile alone is achieved by SyclKittens on the measured single-GPU configurations.Source: §6 · p. 9

6.1 Framework Composition on One GPU

  • 1.22× prefill, 1.65× decode, and 1.59× full-request throughput are the respective geometric-mean gains over torch.compile across the 21 configurations where both systems fit.Source: §6.1 · p. 10
  • Prefill gains grow with context as the FlashAttention-style kernel avoids quadratic score materialization; decode gains are attributed to fewer launches around attention and the projection path.Source: §6.1 · p. 10

6.2 Node-Scale Composition

  • At 24 KiB, SyclKittens beats oneCCL on all four collectives by 1.81–3.31×, while at 96 MiB oneCCL leads all-reduce and all-gather.Source: §6.2 · p. 11
  • Copy-engine overlap beats the blocking PyTorch sendrecv ring by 1.21–2.46× across measured contexts; all-gather exceeds the sequential ring when the gathered KV cache fits in HBM.Source: §6.2 · p. 11

6.3 Multi-GPU Decode

  • At matched GPU counts, SyclKittens outperforms the oneCCL-based decode path at every context, reaching 2.91× at 32K on four GPUs.Source: Figure 7 · p. 11, §6.3 · p. 12
  • At 256K, TP4×CP3 serves requests at 73.3 ms/token using all twelve GPUs; this is a capacity result, not a matched-resource speedup.Source: §6.3 · p. 12

7 Conclusion

The conclusion presents SyclKittens as an Intel GPU tile model whose reusable operations support agent-written kernels and co-designed inference kernels. The accompanying appendices describe measurement protocols, hardware topology, and interface choices.

7 Conclusion

  • SyclKittens provides reusable matrix operand layouts, cooperative cache prefetch, and topology-aware communication, and the paper reports gains over raw SYCL, torch.compile, and a oneCCL-based decode path.Source: §7 · p. 12
  • Whether the interface reduces expert development effort remains untested; the authors identify a human-expert comparison as future work.Source: §7 · p. 12
  • The authors conclude that accelerator interfaces should encode architectural methods while leaving agents’ search focused on workload-specific policies.Source: §7 · p. 12

Appendix overview

  • The appendix overview states that experiments identify their baseline and measurement protocol, with section-specific correctness checks and measurement conventions.Source: Appendix overview · pp. 13–14

A Hardware and Programming-Interface Map

A.1 Measured Fabric Behavior

  • The fabric appendix distinguishes peer-link bandwidth, two-stack all-reduce bus bandwidth, and end-to-end distributed-GEMM speedups; these are not interchangeable estimates of physical link rate.Source: §A.1 · p. 14, Table 5 · p. 14

A.2 Intel Programming-Interface Taxonomy

  • The interface taxonomy contrasts APIs that expose varying schedule structure; SyclKittens lowers tile operations to Intel SPIR-V while kernels choose prefetch ownership and distance.Source: §A.2 · p. 14

B.1 Realization Details

The Intel realization adapts tile operations, operand movement, and correctness fixes to the GPU’s execution model. Controlled agent studies then compare programming contexts under fixed tasks and feedback, including warm refinement of single-GPU GEMM kernels.

B.1 Realization Details

  • Restoring the VNNI transform load after correcting operand marshalling raises single-GPU 8192^3 GEMM throughput from 106 to 147 TFLOP/s, a 39% gain over software repacking.The 2-D block-load path also splits a BlockCount=2 request into two BlockCount=1 loads to stay within the 64-byte hardware load width.Source: §B.1.2 · p. 15
  • The tile-language additions include split-D load/store operations, partial-tile causal masking, and multiplication and minimum reductions beyond the upstream sum and maximum reductions.The measured GEMM path also uses hardware packing for the column-layout B operand to avoid software packing shuffles.Source: §B.1.2 · p. 15
  • The initial Intel port passed 570 of 1,032 unit tests; six distinct fixes account for all 462 failures, bringing the measured-kernel library to 1,032/1,032.Failures included NVIDIA warp-32 and CUDA-shuffle assumptions that do not fit Intel Max GPU’s SIMD16 model, as well as layout, packing, conversion, and reduction issues.Source: §B.1 · p. 14, §B.1.1 · p. 14, Table 8 · p. 16
  • SyclKittens separates L1 cache warming from register loads; cooperative prefetch assigns disjoint operand regions to subgroups, while the measured operand path avoids SLM staging.The Intel Max GPU lacks an SLM matrix-load equivalent; an XMX B-operand requires about 64 SLM reads versus one LoadTransformINTEL from HBM, and measured SLM attention variants ran far slower.Source: §B.1.3 · p. 16, §B.1.4 · p. 16

C.1 Shared Controls

  • The controlled study compares three programming contexts using fixed task inputs, evaluator, and toolchain, with independent starts and at most 20 evaluator calls for cold-iterative agents.The study includes three coding models and single-GPU GEMM, GQA, and distributed GEMM tasks; single-shot agents receive no compilation or execution feedback.Source: C Controlled Coding-Agent Studies · p. 16, §C.1 · p. 17

C.2 Single-GPU Task

  • Single-GPU trials use 4096^3 BF16 GEMM with FP32 accumulation and normalize throughput to a fixed 197.5-TFLOP/s oneDNN baseline.Final correctness requires passing all five held-out seeds; only correct trials receive performance scores.Source: §C.2 · p. 17

C.3 Single-GPU Complete Controlled Results

  • Table 10 gives the ranges omitted from Table 3 in the main paper; a dash means no trial passed all five held-out seeds.Source: §C.3 · p. 17

C.4 Single-GPU Warm-Iterative Refinement

  • Warm feedback improves every eligible Opus GEMM source; the two SyclKittens contexts reach nearly tied frontiers, while examples raise the final median from 70.46% to 86.51%.The Raw source rises from 6.47% to 14.24%; all seven paired sources remain correct and improve.Source: §C.4 · p. 17

C.5 Distributed All-Gather GEMM Studies

The distributed studies examine agent-generated all-gather GEMM kernels and related single-GPU mechanisms, while documenting the evaluation scope and co-design workflow. Results vary by model and context, and the section also reports GEMM-suite and attention performance.

C.5 Distributed All-Gather GEMM Studies

  • Examples raise the distributed single-shot Opus frontier from 36.7% to 59.4% and Sonnet’s median from 0.3% to 34.2%, while Haiku remains below 1%; no generated kernel exceeds the hand baseline or tuned oneCCL path.The distributed task requires a complete 12-GPU row-sharded BF16 all-gather GEMM; examples provide vendor GEMM kernels and a two-GPU all-reduce, but no all-gather GEMM schedule.Source: §C.5 · pp. 17–18
  • All 13 eligible distributed warm-refinement sources remain correct and improve; median paired gains are 5.4 and 5.5 percentage points in the two SyclKittens contexts.One SyclKittens Opus source reaches 106.6% of hand by computing each GPU’s output-row shard before topology-aware output gathering; the unequal source populations make this per-source refinement, not a balanced context comparison.Source: §C.5 · p. 19

C.6 Single-GPU Mechanism Inspection

  • The strongest single-shot Opus distributed kernel uses a 32×32 register base shape, VNNI-compatible B layout, prefetch distance three, 2-D loads, and mma_AB, reaching 81.2% of oneDNN.Raw cold frontiers are 57.5% for Opus, 5.7% for Sonnet, and 1.6% for Haiku.Source: §C.6 · p. 20

C.7 Three-Model Grouped-Query Attention

  • GQA context effects reverse by model: examples lead for Opus at 79.5%, while Raw leads for Sonnet at 18.6%; Haiku stays between 0.7% and 1.3%.All 27 kernels are correct, and the best Sonnet cell falls below the weakest Opus cell at 41.6%; the endpoint-seeded warm campaign covers only six eligible Opus sources, which all improve.Source: §C.7 · pp. 20–21

C.8 Data Quality and Protocol Deviations

  • The single-shot and cold-iterative comparisons retain their planned balance despite protocol deviations; warm-iterative analyses remain paired because eligibility determines their source populations..Source: §C.8 · p. 22

C.9 AI Use in SyclKittens development

  • Engineers and coding agents jointly developed and tuned the suite, while engineers owned experiments and acceptance; generated code entered only after correctness and performance checks and human review.The controlled autonomous study is distinct from this engineer-agent suite.Source: §C.9 · p. 22

D BF16 GEMM

  • The BF16 SyclKittens GEMM suite achieves 0.963× oneDNN in geometric-mean throughput across 27 shapes.Measurements use BF16 output, shape-adaptive launch policies, within-trial co-measurement, and six-node aggregation.Source: D BF16 GEMM · p. 22

D.1 Configuration

  • GEMM FLOP counts are 2MKN for both implementations, so their TFLOPS figures are directly comparable.Source: §D.1 · p. 22

D.2 Kernel Implementation and Key Techniques

The kernel uses register-resident DPAS tiles, transformed block loads, and cooperative prefetch ownership, with shape-specific tuning. The section reports gains from de-duplication and lookahead, correctness checks, comparisons with Triton-XPU and oneDNN, and attention results under specified configurations.

D.2 Kernel Implementation and Key Techniques

  • Each subgroup streams K in 32-element chunks, loads row-major A and column-major B tiles through 2-D block loads, applies the VNNI transform to B, and contracts with DPAS.Collaborative partial-tile prefetch assigns disjoint panel regions so each shared operand byte is warmed once per K-step; the workgroup warms 32,768 bytes per step versus about 118,784 bytes for masked full-tile prefetch.Source: §D.2 · p. 22, §D.2.1 · pp. 23–24

D.3 Methodology

  • Measurements interleave oneDNN and SyclKittens, use 10 warmups and 30 measured iterations, and aggregate per-shape ratios through node medians.Outputs are checked against torch.mm by relative Frobenius norm; all measured configurations pass, with maximum error 1.663 × 10^-3.Source: §D.3 · p. 24

D.4 De-Duplicated-Prefetch Performance

  • 1.056× geometric-mean gain from de-duplication and shape-specific prefetch distances yields the final 96.3% of oneDNN configuration across 27 shapes.The cumulative tail compares snake traversal, partial-tile ownership de-duplication, and shape-specific prefetch distance.Source: §D.4 · p. 24, Table 21 · p. 24

D.5 Correctness

  • All 27 shapes pass the relative-Frobenius gate below 0.02, with uniform relative error of 1.66 × 10^-3 and no per-shape outlier.The BF16-output store retains the FP32 output path’s accuracy.Source: §D.5 · p. 24

D.6 Triton-XPU Comparison

  • 1.087× geometric-mean and 1.030× median throughput over Triton-XPU 3.6; SyclKittens’ cross-node median exceeds Triton on all 27 shapes.The comparison uses independently autotuned Triton configurations and reports ratios reduced by medians within each node and then across six nodes.Source: §D.6 · p. 24, Table 24 · p. 25

D.7 Split-K / Stream-K Design

  • Split-K runs 2.4–6.9× faster than the default kernel and exceeds oneDNN by up to 1.21× for K≤16384; at K=32768, oneDNN is faster.The path targets occupancy-starved shapes with small M and N, large K, and an L2-resident output; its atomic-add epilogue is bit-exact in all 10 unit tests.Source: §D.7 · p. 25, §D.7.2 · p. 25

E Attention Forward (MHA and GQA)

  • The fence-free adaptive forward-attention kernel reaches 71.7 TFLOPS for MHA and 75.7 TFLOPS for GQA; speedups over IPEX SDPA are 1.09–1.21× for MHA and 1.70–4.0× for GQA.MHA is measured at batch four, while the GQA performance tables measure batch-one full-sequence queries.Source: E Attention Forward (MHA and GQA) · p. 25, §E.1 · p. 25

E.1 Configuration

  • GQA suite comparisons use IPEX SDPA with enable_gqa=True and its grouped-head fallback, which repeats KV heads; the controlled agent study instead uses native PyTorch XPU SDPA without IPEX.No Triton grouped-attention baseline is reported.Source: §E.1 · p. 25

E.2 Kernel Implementation and Key Techniques

This section describes SyclKittens attention, normalization, and RoPE kernels, their hardware-aware scheduling and data movement, and their measured performance. It also reports that the RoPE implementation is bit-exact against an fp32 reference at all six tested shapes, matching torch.compile.

E.2 Kernel Implementation and Key Techniques

  • The single-pass, register-resident FlashAttention-style kernel streams KV tiles, keeps online-softmax state in registers, and prefetches upcoming K/V loads while DPAS consumes the current tile.Avoiding forced root synchronization improved large-N throughput by up to ∼16×; split-K improved performance by 0.6–2.9% at N ≥ 4096.Source: §E.2 · pp. 25–26

E.4 Delivered Performance

  • Non-causal MHA reached 1.09× over SDPA and causal MHA 1.21×; non-causal GQA reached 1.70× and causal GQA 4.0× over SDPA.At N = 512, the selected causal MHA variant reached 1.42× IPEX and 1.87× Triton.Source: Table 27 · p. 26, Table 28 · p. 26, §E.4 · p. 26

E.5 Autotune Grid and Selection Policy

  • GQA dispatch depends on both sequence length and batch size: at causal N = 1024, reusing the B = 1 selection at B = 4 took 681 µs versus 569 µs with 16 query rows.The 2-D policy shifts the seq_q = 16 crossover to smaller N as batch fills the machine; the B = 1 selection was 20% slower at B = 4 and the difference was 38% at B = 8.Source: §E.5 · p. 27

E.6 Correctness

  • All MHA configurations pass correctness checks, including causal seq_q = 8 at ≤0.011; all GQA configurations pass tolerance 0.10, with causal seq_q = 8 at ≤0.0134.The RMSNorm and LayerNorm configurations also pass, and the residual store is exact.Source: §E.6 · pp. 27–28, §F.6 · p. 30

F.2 Kernel Implementation and Key Techniques

  • Eight-subgroup row decomposition and 128-bit coalesced load_wide/store_wide transactions raise measured LayerNorm bandwidth from ∼570 to ∼970–1100 GB/s.The implementation also fuses residual addition, normalization, and affine operations in one launch.Source: §F.2 · p. 29

F.4.1 RMSNorm (𝐷= 2048). Plain, geomean 1.12× over torch.compile, 15.6× over eager.

  • Fused residual-add and norm had a geomean 1.04× over torch.compile, while SyclKittens held a warm median ∼24 µs dispatch floor across the sweep versus torch.compile at ∼55–61 µs.Source: §F.4.1 · p. 29, §F.5 · p. 29

F.4.2 LayerNorm (𝐷= 4096). Plain, geomean 1.42× over torch.compile, 2.78× over eager.

  • Fused residual-plus-norm LayerNorm (4×) achieved a 1.12× geometric-mean speedup over torch.compile; on HBM-bound shapes (B≥16), SyclKittens sustained 623–997 GB/s on LayerNorm, or 1.08–1.63× torch.compile.At the two large RMSNorm cells, both paths approached the ∼1.3 TB/s per-GPU roofline and differed by 0.99–1.05×.Source: §F.4.2 · p. 29

F.7 Rotary Position Embedding (RoPE) with 128-bit Chunk-Pair Mapping

  • RoPE reached 1.09× over torch.compile for cold bandwidth-bound shapes and 1.98× for warm decode, with 1020 GB/s at B = 64.Its 128-bit chunk-pair mapping gives each work-item a rotary pair locally, requiring no SLM, barrier, or shuffle; warm decode latency stayed near-flat at ∼27–31 µs.Source: §F.7 · p. 30, §F.7.2 · p. 30, §F.7.3 · p. 31
  • Relative error was 0.0 at all six shapes, matching torch.compile.Every configuration was validated against an fp32 reference with a relative-error tolerance of 2e-2.Source: §F.7.4 · p. 31
  • The halved rotary is computed in fp32 and rounded to bf16 identically to the reference, producing a bit-exact result.Source: §F.7.4 · p. 31
  • Every configuration is validated against an fp32 reference before timing, using a relative-error tolerance of 2e-2.Source: §F.7.4 · p. 31

F.8 Single-Token Decode GEMV (𝑀= 1)

SyclKittens’ single-token GEMV uses split-K to increase parallel work and reduce decode latency against oneMKL. The section also introduces fused Q/K rotary as a way to combine launches and reuse shared cosine/sine data.

F.8 Single-Token Decode GEMV (𝑀= 1)

  • The five decode-time weight projections account for approximately 64% of per-token decode latency and are limited by bf16 weight-streaming bandwidth.Each processes one token (M=1) and streams its weight matrix once.Source: §F.8 · p. 31
  • Split-K assigns multiple subgroups to each output row’s K reduction, increasing parallel work as output-row count shrinks.The kernel uses 128-bit coalesced loads, software-prefetched chunks, fp32 accumulation, and direct bf16 output; the stated configuration requires K divisible by 2048.Source: §F.8.2 · pp. 31–32
  • 1.36× cold-HBM geomean over oneMKL; summed per-token GEMM latency falls from 36.3 ms to 27.2 ms (1.33×).Across cold and L2-resident regimes, SyclKittens/oneMKL ratios are 1.24–1.53× and 1.02–1.90×, respectively.Source: §F.8.3 · p. 32
  • SyclKittens’ maximum absolute error matches oneMKL across the five shapes, with results matching oneMKL to the last bf16 ULP band.Both paths accumulate in fp32 and round once to bf16.Source: §F.8.5 · p. 33

G.1 Fused Q/K Rotary

  • Fusing Q/K rotary combines two launches into one and permits cached reuse of shared cosine/sine data across Q and K work-items.The comparison uses eager PyTorch and torch.compile on Intel Max GPU; the fused path retains the same halved-rotary math.Source: G Op-Fusion Microbenchmarks · p. 33, §G.1 · p. 33

G.1.2 Kernel implementation and key techniques.

The fused rotary kernel concatenates Q and K work-items into one launch and shares a per-row worker. The section also reports fused-rotary performance and a topology-aware distributed GEMM path.

G.1.2 Kernel implementation and key techniques.

  • One concatenated Q/K id space routes each work-item to the appropriate tensor while both branches call the same per-row rotary worker.Each work-item owns one eight-element chunk-pair; with D=128, there are eight chunk-pairs per row, and the kernel uses coalesced loads and stores without SLM, barriers, or shuffles.Source: §G.1.2 · pp. 33–34
  • The rotary measurements use interleaved round-robin timing against eager PyTorch and torch.compile, with warm L2-resident and cold L2-flushed regimes.The methodology checks rotated Q and K against an fp32 reference before timing.Source: §G.1.3 · p. 34, §G.1.4 · p. 34

G.1.4 Delivered

  • 1.41× over torch.compile for cold bandwidth-bound shapes and 2.06× for warm decode shapes in fused Q/K rotary.Cold-shape geomean is 9.03× over eager; warm-shape geomean is 17.2× over eager. The fused path is 1.26–1.69× faster cold and 1.98–2.13× faster warm, with 138–149 µs warm latency.Source: Table 46 · p. 34, Table 47 · p. 34, §G.1.4 · p. 34

H Distributed GEMM

  • On 12 GPUs, the topology-aware input-gather path reaches 84% and 93% of its co-measured GEMM-only baseline at 61443 and 122883.It all-gathers A, replicates each source stripe over MDFI to the partner stack, then computes one local GEMM.Source: H Distributed GEMM · p. 35

H.1 Configuration

  • The distributed GEMM configuration uses six GPU packages with two stacks each, connected within packages by MDFI and across packages by XeLink.The operation starts with row-sharded A and replicated B and returns full replicated C; the hand GEMM core uses bf16 register-resident 32×64×32 DPAS tiling with snake swizzle.Source: §H.1 · p. 35

H.2 Kernel Implementation and Key Techniques

The distributed-GEMM path reduces cross-package traffic by routing each shard once per remote package, while materializing inputs before GEMM avoids stalls from remote operand reads. The section also describes engineer-agent refinement and context-parallel attention’s ring and all-gather approaches.

H.2 Kernel Implementation and Key Techniques

  • 1.42×–2.15× end-to-end distributed-GEMM speedup over the flat gather pipeline comes from sending each stripe across XeLink once, then replicating it over MDFI.Materialize-then-GEMM is faster than interleaving gathers or reading remote A inside the kernel, which stalls the DPAS pipeline.Source: §H.2 · p. 35

H.3 Methodology

  • The distributed-GEMM measurements run all 12 GPUs concurrently in one process and compare methods within the same run against the GEMM-only baseline.Each method uses the median of 30 timed iterations after 10 warmups and a 2-second turbo burn-in.Source: §H.3 · p. 36

H.4 Engineer-Agent Co-Design

  • The engineer-agent refinement matches the incumbent: its median ratio is 1.007×, with a 95% confidence interval of 0.93× to 1.04×.The refinement was certified through the promotion protocol, but its margin is small relative to run-to-run variation.Source: §H.4 · p. 36
  • The refinement protocol retains the strongest validated agent source as an incumbent and promotes an engineered variant only after it wins the same correctness and timing protocol.The proposed search sequence enumerates collective placement, carries topology routing across that placement, then retunes for the resulting local shard geometry.Source: §H.4 · p. 36

H.5 Topology-Aware All-Gather with Collaborative Prefetch

  • The production topology-aware path reaches 93% of the local compute ceiling on the larger shape and retains 84% on the smaller, communication-bound shape.Both shapes pass the analytic constant-fill correctness gate; these results use random-normal inputs and the cooperative-prefetch build.Source: §H.5 · p. 36, Table 49 · p. 36, §H.6 · p. 37

I Context-Parallel Attention (Ring and All-Gather)

  • Context-parallel attention combines ring and all-gather KV distribution, measured on three clean nodes with both paths checked against the same correctness gate.The harness pairs PyTorch communication paths with Triton and IPEX/aten flash kernels.Source: I Context-Parallel Attention (Ring and… · p. 37, §I.3 · p. 38

I.2 Kernel Implementation and Key Techniques

  • The ring rotates KV shards between GPUs and merges online-softmax state, while all-gather collects the full context once for a single flash pass.Ring uses O(N_per) memory per GPU; all-gather uses O(N_total) and runs out of memory when the gathered context exceeds HBM.Source: I Context-Parallel Attention (Ring and… · p. 37, §I.2 · p. 38
  • The ring’s per-step sequence computes partial attention, merges the online-softmax state, then copies the next KV shard into the alternate local-HBM buffer.The in-order queue serializes compute and transfer; a later variant uses copy engines to overlap rotation with compute.Source: §I.2 · pp. 37–38

I.4 Delivered Performance

Context-parallel attention trades memory footprint against performance: all-gather is faster when the full context fits in HBM, while ring supports longer contexts. The section also reports ring and all-gather comparisons, correctness checks, and multi-GPU collective results.

I.4 Delivered Performance

  • 1.85× to 1.05× against sendrecv/aten as N_per grows from 256 to 16384, while the ring remains 1.76–2.89× faster than sendrecv/triton.The sequential schedule is reported for a B=1 context sweep across six GPU packages.Source: §I.4.1 · p. 38, Table 52 · p. 39
  • 1.21–2.46× the aten ring throughput, overlapping shard rotation with compute raises the six measured ring configurations to 1.7–69.0 TFLOP/s per GPU.At N_per from 256 to 8192, the flash-core share rises from 20% to 65%, while reduction falls from 41% to 6% and rotation from 40% to 29%.Source: §I.4.1 · p. 38, Table 52 · p. 39, Table 53 · p. 39
  • 4.26× at N_per=256 narrowing to 1.17× at 16384, all-gather’s advantage over the sequential ring declines as the shard grows.All-gather remains 1.65–2.79× faster than allgather/triton; the tuned oneCCL-plus-IPEX path overtakes it only at the two largest shards.Source: §I.4.2 · p. 38, Table 54 · p. 39
  • 1.17–4.26× faster than the sequential ring across the measured B=1 sweep, all-gather pays the inter-package cost once when the gathered context fits in HBM.All-gather uses O(N_total) memory per GPU and runs out of memory beyond HBM; ring stays O(N_per), supporting longer contexts.Source: Figure 10 · p. 39, §I.4.2 · p. 40

I.5 Correctness

  • All functional configurations pass the 0.05 tolerance; the all-gather path is also checked in-binary against a full-context online-softmax reference.At N_per=256, that all-gather reference check has relative error 0.0012.Source: §I.5 · p. 40

J Multi-GPU Collectives

  • All four collectives win on every node at 24 KiB, while at 96 MiB oneCCL leads all-gather and all-reduce and SyclKittens all-to-all wins on all six nodes.Reduce-scatter favors SyclKittens on four of six nodes at 96 MiB; medium all-gather, reduce-scatter, and all-reduce results vary by node.Source: J Multi-GPU Collectives · p. 40

J.1 Configuration

  • The 12-GPU collective comparison uses one SyclKittens process with shared-context peer pointers against one oneCCL process per GPU.SyclKittens dispatches each collective adaptively between latency-oriented and topology-aware bandwidth paths.Source: §J.1 · p. 40

J.2 Kernel Implementation and Key Techniques

SyclKittens collectives use peer access and topology-aware paths, with performance varying by collective and payload. The evaluated kernels also support end-to-end Llama-3.1-8B gains when composed with TorchInductor-generated regions.

J.2 Kernel Implementation and Key Techniques

  • The collective kernels use direct peer read/write and adapt between latency- and bandwidth-oriented paths; techniques include fan-out writes, hierarchical reductions, and destination-split all-to-all copies.Hierarchical reduction sums the two stacks within each package over MDFI before reducing across the six packages over XeLink.Source: §J.2 · p. 40

J.3 Methodology

  • Collective measurements pair SyclKittens and oneCCL on the same 12 GPUs and links, with per-node paired speedups computed from bracketing medians.SyclKittens points are full-buffer verified against reference results, with maximum absolute error within bf16 rounding.Source: §J.3 · p. 41

J.4 Delivered Performance

  • At 24 KiB all four collectives win on all six nodes; at 96 MiB oneCCL leads all-gather and all-reduce, while SyclKittens all-to-all wins on all six nodes.Medium all-gather, reduce-scatter, and all-reduce cells vary by node; the six-node figure reports paired same-node brackets.Source: §J.4 · p. 41, Figure 11 · p. 42
  • 1.1–1.54× multi-GPU decode TPOT improvement follows replacing 64–96 per-token host waits with a system-scope atomic peer-counter barrier.The synchronization change does not alter arithmetic and passes the same correctness checks.Source: §J.4 · p. 41

J.5 Correctness

  • All SyclKittens collectives pass full-buffer verification against reference sums, gathers, or scatters, with maximum absolute error at most 0.01.The oneCCL baseline uses the torch.distributed collective.Source: §J.5 · p. 41

K End-to-End Llama-3.1-8B

  • 1.59× geometric-mean full-request throughput over torch.compile alone on the 21 configurations both paths run, composing SyclKittens operations with TorchInductor-generated regions.The composed path also serves the complete N=8192 grid, including cells where the baselines run out of memory.Source: K End-to-End Llama-3.1-8B · p. 42

K.1 Evaluated Kernel Set

  • Each end-to-end kernel selection is the fastest correct implementation in its operation-specific study, and each run records which implementation executes.The torch.compile baseline fuses pointwise and reduction regions but graph-breaks at SDPA on this backend.Source: §K.1 · p. 42, §K.2 · p. 42, Table 56 · p. 42

K.3 Operations across the SDPA Graph Break

  • The end-to-end path places SyclKittens GQA flash attention, joint Q/K rotary, concatenated QKV projection, and residual-add plus RMSNorm across the SDPA graph break.Inductor retains RMSNorm, SwiGLU, and the remaining pointwise and epilogue regions.Source: §K.3 · p. 42

K.4 End-to-End Composition with TorchInductor

The evaluation composes SyclKittens operations with TorchInductor across the SDPA graph break, then compares multi-GPU decode with oneCCL and explores TP×CP layouts. Results include faster matched-GPU decode and context-dependent long-context layouts.

K.4 End-to-End Composition with TorchInductor

  • SyclKittens replaces regions spanning the SDPA graph break, while Inductor retains the remaining pointwise, reduction, and epilogue regions.Stand-alone norm and SiLU stay with Inductor because it matches or exceeds the corresponding SyclKittens kernels; both runs use identical compilation settings, differing only in four SyclKittens operations.Source: §K.4 · p. 42, §K.4.1 · p. 43
  • At prompt length 256, prefill ratios are 0.91×–1.03×; from 512 onward they rise with context, reaching 2.03× at 1 × 8192.At 256, small projection GEMMs do not amortize fused QKV, rotary, and residual launches; from 512, the flash-attention path replaces baseline attention across the graph break.Source: §K.4.2 · p. 43
  • 1.65× geometric-mean decode improvement, reaching 1.96× at 1 × 256, accompanies fewer launches in the composed path.The SDPA graph break prevents Inductor from fusing the per-step attention prologue and epilogue.Source: §K.4.2 · p. 43
  • At N=8192, torch.compile runs out of memory at B≥2 and eager at B≥4, while the composed path serves BS 1–8.The passage attributes the composed path’s coverage to flash attention’s smaller memory footprint.Source: §K.4.2 · p. 43
  • 1.56× end-to-end geometric mean: the composed path outperforms Inductor alone across the N≤2048 subset.The gain combines Inductor’s pointwise/norm fusion with cross-op and attention fusions spanning the SDPA graph break.Source: §K.4.2 · p. 43

K.5 Distributed Multi-GPU Decode (Single-Process TP+CP vs. oneCCL)

  • 1.52–2.91× faster: SyclKittens beats oneCCL decode on matched two- and four-GPU configurations across contexts where both run.The reported TPOT speedups cover the contexts where both implementations run.Source: §K.5.4 · p. 44, Table 61 · p. 44
  • All-reduces account for 3.4–3.7% at short context and less than 1.5% after the KV scan dominates; two-GPU oneCCL TPOT rises from 43 to 105 ms from 1K to 32K, versus 23 to 36 ms for SyclKittens.The collective percentages come from an instrumented profile, not the final overlapped TPOT table.Source: §K.5.5 · p. 44
  • At 64K and 128K, six-GPU TP2×CP3 is 1.76× and 2.11× faster than flat P=4; TP4×CP3 serves 256K at 73.3 ms/token.Table 62 reports end-to-end TP×CP decode TPOT across the node’s configurations.Source: §K.5.7 · p. 45, Table 62 · p. 45

Related papers

Loading 2610.04277v1…