Tasks/Systems & Hardware

Paged Ragged GQA Decode Kernel Optimization

Accelerate grouped-query attention over paged, variable-length KV caches and unseen shapes

Systems & HardwareLLM inferenceattention kernels
Background

Attention kernels for large-language-model serving read a key-value cache split into fixed-size pages, and making this memory-bound decode step fast remains an open engineering problem. The starting point is a plain vectorized implementation that gathers each request's pages and computes grouped-query attention, in which many query heads share one cached key-value head. The work is to redesign that computation for far lower latency without changing what it returns. Speed must survive shuffled page tables, wildly uneven sequence lengths, and shapes withheld until scoring.

instruction.mdthis is what the agent is given

You inherit a correct batch-vectorized PyTorch implementation of grouped-query decode attention over a paged, variable-length KV cache. Make it as fast as possible while it keeps computing the same thing within the accuracy the verifier enforces.

Hard Constraints

  • Keep the callable signature in methods/main/solver.py unchanged. Modify implementation files only under methods/main.
  • Inputs are CUDA tensors. q has shape (batch, query_heads, head_dim). k_cache and v_cache have NHD layout (physical_pages, page_size, kv_heads, head_dim).
  • page_indptr, page_indices, and last_page_len are CUDA int32 tensors. They encode a CSR page table with at least one page per request; the final valid page length is in [1, page_size].
  • Query and cache tensors share dtype, either float16 or bfloat16. query_heads is divisible by kv_heads; each contiguous query-head group maps to one KV head.
  • Return one fresh CUDA tensor with shape (batch, query_heads, head_dim) and the input floating dtype. Do not mutate or alias any input.
  • Implement the documented attention for every request: gather only pages named by its page-table row, trim the final page, optionally apply the specified Llama half-split RoPE to query position kv_len - 1 and key positions 0..kv_len-1, retain only the last window_left + 1 tokens when window_left >= 0, form scores = q @ k.T * sm_scale, apply scores = logits_soft_cap * tanh(scores / logits_soft_cap), apply softmax in each GQA head, and multiply by values. RoPE frequencies are theta^(-arange(0, head_dim, 2)/head_dim) / rope_scale; mode 0 means no positional encoding and mode 1 means ROPE_LLAMA.
  • The sealed panel uses page tables, request lengths, shapes, dtypes and tensor values you have not seen. Your implementation must stay correct and fast on them. Do not enumerate visible cases, fixed seeds, or any other property you can only observe in the public panel.
  • Caching an output across calls, mutating metadata, or replacing the timing primitives is detected and scored as incorrect.
  • Submission code may import torch, triton, standard typing/math helpers, and sibling modules. It may not access files, processes, clocks, environment secrets, network services, verifier paths, or runtime introspection.

What You Have

  • methods/main/solver.py: the correct batch-vectorized PyTorch starting implementation covering the full callable ABI.
  • problems/visible_spec.json: a public development panel with real shuffled physical page tables and ragged request lengths.
  • protocol.py: the public data contract, the reference semantics, and the exact measurement protocol the verifier uses.
  • selfcheck.py: an isolated visible evaluator. It generates fresh values, checks every output against an independent float32 formulation, and reports visible_geomean_speedup against the inherited starter. Higher is better.
  • last_page_len gives each request's valid tail length; slots past it are unwritten cache memory and must not be attended to.
  • /app/methods/experiment_log.md: the append-only experiment ledger. Record each attempt's visible latency, its visible_geomean_speedup, and the keep/revert decision, and save the matching source under /app/methods/versions/vN/.

What You Submit

The contents of /app/methods, with your final implementation at /app/methods/main/solver.py. The evaluator imports:

from solver import paged_gqa_decode

Module-level compilation caches and reusable workspaces are allowed. They must be keyed by general runtime properties and stay correct when a fresh process presents unseen cases. Planning and compilation happen during untimed warmup; timed calls still receive fresh q, k_cache, and v_cache tensors.

How It Is Judged

Each sealed case runs the verifier's own hash-bound copy of the inherited starter immediately before your candidate and immediately after it. All three runs use fresh isolated processes, identical inputs and seeds, the same correctness checks, and the same timer. The case score is the bracketed baseline latency divided by your candidate latency, and the cases are combined by geometric mean. The inherited starter therefore scores 1.0x.

Before a case scores at all it must pass output structure, finite values, numerical accuracy, input immutability, no-alias, fresh-call, timing-integrity, dispersion, and baseline-drift checks. Any failed case zeroes the entire score. Timing measurement is defined in protocol.py; input generation, validation and output transfer are all outside the timed window.

Run the visible loop with:

python /app/selfcheck.py

Rollouts

$56.04Spend
115.2MTokens
68Versions, 42 kept

On the visible set

0 5 10 15 20 25 0 15 30 45 60 Agent step Visible-panel speedup vs verifie ↑ v0 v1 v2 v3 v4 v5 v6 v7 v8 v9 v10 v11 v12 v13 v14 v15 v16 v17 v18 v19 v20 v21 v22 v23 v24 v25 v26 v27 v28 v29 v30 v31 v32 v33 v34 v35 v36 v37 v38 v39 v40 v41 v42 v43 v44 v45 v46 v47 v48 v49 v50 v51 v52 v53 v54 v55 v56 v57 v58 v59 v60 v61 v62 v63 v64 v65 v66 v67
keptrolled backsubmitted
  1. v0The agent started from the inherited vectorized PyTorch decode baseline.1.00490 min · $0.00
  2. v1The agent fused decode into one Triton pass over the paged cache.2.1964
  3. v2The agent tried fast trig approximations, and the gain never showed.2.1956
  4. v3The agent precomputed rotation tables to stop the kernel spilling.3.8727
  5. v4The agent tried a relative-position rotation trick, and heavy cases lost.3.6887
  6. v5The agent split long histories into planned chunks to fill the machine.16.2376
  7. v6The agent cut chunk length so the long tail got more blocks.16.7690
  8. v7The agent tried deriving chunks from a block-count target, and odd sizes hurt.16.6039
  9. v8The agent shrank chunks further, and per-chunk overhead took over.15.1089
  10. v9The agent doubled the inner token tile to halve loop iterations.22.2495
  11. v10The agent pushed that tile to 128 everywhere, and rotation spilled.19.2911
  12. v11The agent sized the token tile by whether rotation was on.22.3361
  13. v12The agent rewrote the soft cap as an exact sigmoid identity.22.5256
  14. v13The agent stored the split partials in low precision to halve traffic.22.7401
  15. v14The agent reduced the split bookkeeping to one log-sum-exp number.22.9636
  16. v15The agent tried fewer warps in the reduction, and it came out tied.22.9515
  17. v16The agent added a pipeline stage and repeated the run to be sure.23.0179
  18. v17The agent added another stage, and the gain did not repeat.22.9427
  19. v18The agent tried one wide rotation load, and the reshaping cost more.19.8840
  20. v19The agent precomputed the page lookup, and plain arithmetic was cheaper.21.2479
  21. v20The agent tried rotating keys once per call instead of inside the loop.22.1426
  22. v21The agent tuned that pre-rotation pass on its own branch.22.2332
  23. v22The agent widened the pre-rotation tile, and throughput fell.22.1387
  24. v23The agent routed each request to inline or pre-rotation by its workload.23.7886
  25. v24The agent widened the query rotation tile, and nothing improved.23.7169
  26. v25The agent gave the key rotation more warps, and it slowed.23.6865
  27. v26The agent starved the query rotation of warps, and it slowed too.23.6540
  28. v27The agent specialized tile width to each route's register budget.23.8991
  29. v28The agent widened that tile to 256, and spills took the gain back.23.5783
  30. v29The agent cut warps everywhere, and the inline rotation path collapsed.22.2662
  31. v30The agent chose warp counts per route instead of one global setting.24.2566
  32. v31The agent retuned the plain path's tile until its spills disappeared.24.6136
  33. v32The agent made every simplified route spill-aware.24.8285
  34. v33The agent folded constant scaling onto the host side.24.9122
  35. v34The agent tidied the planner and named its routing crossover.24.8961
  36. v35The agent widened the head tile, and the affected cases went flat.24.7551
  37. v36The agent tried a fast path for single-split requests, and it did not hold.24.8357
  38. v37The agent revisited the relative-position rotation trick, and it still lost.24.8042
  39. v38The agent scaled the token tile to the head dimension for unseen shapes.24.9219
  40. v39The agent refined that rule for the pre-rotated route at one head size.24.8716
  41. v40The agent pushed the smallest head dimension to much wider tiles.24.8131
  42. v41The agent kept the wider tiles only on the routes that liked them.24.9503
  43. v42The agent tried simplifying the host dispatch, and it stalled the queue.24.7822
  44. v43The agent matched warp count to tile pressure on the small-head path.24.9987
  45. v44The agent lengthened split chunks for the smallest head dimension.24.9304
  46. v45The agent lengthened them again and the unseen-shape panel agreed.24.8789
  47. v46The agent pushed split length until the routes began to disagree.24.8812
  48. v47The agent tried per-route split sizes, and planner overhead ate the gain.24.8952
  49. v48The agent carried the split-length rule over to the middle head size.24.9309
  50. v49The agent doubled that split length and the panel still improved.24.8567
  51. v50The agent doubled it once more, and the parallelism collapsed.24.9532
  52. v51The agent widened the plain route's tile at the middle head size.24.7900
  53. v52The agent extended that wider tile to the rotated route as well.24.8961
  54. v53The agent repeated the widening on the smallest head dimension.24.8247
  55. v54The agent extended it to the rotated route there too.24.9415
  56. v55The agent doubled the tile again, and the plain route fell apart.24.9883
  57. v56The agent tried rescuing the oversized tile with more warps, and it barely moved.24.8814
  58. v57The agent kept the safe tile and gave it more warps instead.24.8644
  59. v58The agent carried the warp bump to the middle head size.24.8272
  60. v59The agent gated that warp bump on how much parallelism the case had.24.8438
  61. v60The agent gave the pre-rotated route wider tiles and more warps.24.9227
  62. v61The agent transferred that pairing to the middle head size.24.8252
  63. v62The agent widened the key rotation tile again, and every slice slowed.24.7958
  64. v63The agent widened the query rotation tile for the smaller head sizes.24.8400
  65. v64The agent widened it once more and repeated the measurement to be sure.24.8031
  66. v65The agent measured where routing should switch instead of guessing.25.0847
  67. v66The agent made the routing threshold aware of batch-driven wave cliffs.24.9369
  68. v67The agent tried shrinking tiles to the window size, and efficiency won out.24.8301

On the hidden set

Original metricNormalised score
Starter1.00x0.00
Frontier-calibrated reference73.55x0.60
Upper233.38x1.00
This run (GPT-5.6-sol)23.87920.1906
116 minWall clock
$43.25Spend
166.9MTokens
6Versions, 5 kept

On the visible set

0 15 30 45 60 75 0 2 3 4 Agent step Visible-panel speedup vs verifie ↑ v0 v1 v2 v3 v4 v5
keptrolled backsubmitted
  1. v0The agent inherited a vectorized PyTorch paged-GQA decode baseline1$4.38
  2. v1The agent rewrote decode as a Triton flash-decoding kernel resolving paging in-kernel70.96262 min · $8.76
  3. v2The agent rotated keys in input dtype and chained the dots, removing spills72.03179 min · $13.77
  4. v3The agent ordered longest chunks first and added stride guards for robustness71.28394 min · $18.90
  5. v4The agent allocated the output with empty_like off a pinned prototype71.373100 min · $20.35
  6. v5The agent sized tiles from the card's reported shared-memory limit71.364113 min · $24.16

On the hidden set

Original metricNormalised score
Starter1.00x0.00
Frontier-calibrated reference73.55x0.60
Upper233.38x1.00
This run (Opus 5)71.37220.5689
-Wall clock
$4.33Spend
21.7MTokens
5Versions, 4 kept

On the visible set

0 4 8 12 16 20 0 1 2 3 4 Agent step Visible-panel speedup vs verifie ↑ v0 v1 v2 v3 v4
keptrolled backsubmitted
  1. v0The agent inherited a vectorized PyTorch paged-GQA decode baseline1
  2. v1The agent fused decode into one Triton kernel with online softmax1.487
  3. v2The agent adopted two-stage FlashDecoding split-KV with dynamic sequence chunking19.89
  4. v3The agent added adaptive group-size and sequence-chunking heuristics20.377
  5. v4The agent loaded precomputed Q-RoPE directly and tuned block bounds20.528

On the hidden set

Original metricNormalised score
Starter1.00x0.00
Frontier-calibrated reference73.55x0.60
Upper233.38x1.00
This run (Gemini 3.7 Flash)19.77210.1570
-Wall clock
-Spend
-Tokens
11Versions, 7 kept

On the visible set

0 15 30 45 60 0 2 5 8 10 Agent step Visible-panel speedup vs verifie ↑ v0 v1 v2 v3 v4 v5 v6 v7
keptrolled backsubmitted
  1. v0The agent inherited a vectorized PyTorch paged-GQA decode baseline1
  2. v1The agent built a two-kernel Triton split-K flash-decode with cached work plans21.055
  3. v2The agent precomputed fp16 RoPE tables and tuned both kernel stages56.875
  4. v3The agent shrank the KV chunk from 256 to 19260.424
  5. v4The agent added evict_first streaming loads and a one-wave chunk guard60.439
  6. v5The agent moved softmax to the log2 domain with overflow-safe exp2-tanh60.934
  7. v6The agent retuned stage two's block size and pipeline stages60.7
  8. v7The agent folded LOG2E into the soft-cap scalar60.6
  9. v7a-atomicThe agent tried a single-kernel atomic fp32 reduction and reverted itlocal ~42 us/call
  10. v7b-fulldotThe agent tried a full-width QK dot for the no-RoPE pathcase6 26.8->29.7 us
  11. adaptive-chunk-maxctasThe agent tried a max-CTA chunk tie-break rulegeo 29.4 vs 28.8

On the hidden set

Original metricNormalised score
Starter1.00x0.00
Frontier-calibrated reference73.55x0.60
Upper233.38x1.00
This run (Kimi K3)60.20080.4943
-Wall clock
$8.03Spend
12.4MTokens
2Versions, 1 kept

On the visible set

0 4 8 12 16 0 0 0 1 1 Agent step Visible-panel speedup vs verifie ↑ v0 v1
keptrolled backsubmitted
  1. v0The agent inherited a vectorized PyTorch paged-GQA decode baseline1
  2. v1The agent fused a Triton paged GQA decode with online softmax17.714

On the hidden set

Original metricNormalised score
Starter1.00x0.00
Frontier-calibrated reference73.55x0.60
Upper233.38x1.00
This run (Grok 4.6)15.85570.1250
137 minWall clock
$54.92Spend
220.6MTokens
6Versions, 5 kept

On the visible set

0 7.5 15.0 22.5 0 2 3 4 Agent step Visible-panel speedup vs verifie ↑ v0 v1 v2 v3 v4 v5
keptrolled backsubmitted
  1. v0The agent inherited a vectorized PyTorch paged-GQA decode baseline1$9.14
  2. v1The agent wrote a Triton split-token flash-decode with inline RoPE tables18.32691 min · $18.28
  3. v2The agent moved rotation to fp16 and capped registers at 12823.15899 min · $21.97
  4. v3The agent set num_stages to one, cutting spills and shared memory23.686105 min · $23.44
  5. v4The agent gated maxnreg=128 behind a plan-build bitwise self-check23.49120 min · $30.02
  6. v5The agent removed the norm kernel using per-group atomic completion counters26.06133 min · $34.95

On the hidden set

Original metricNormalised score
Starter1.00x0.00
Frontier-calibrated reference73.55x0.60
Upper233.38x1.00
This run (DeepSeek V4 Pro)25.81310.2069
220 minWall clock
$149.66Spend
29.9MTokens
12Versions, 8 kept

On the visible set

0 15 30 45 60 0 2 5 8 10 Agent step Visible-panel speedup vs verifie ↑ v0 v1 v2 v3 v4 v5 v6 v7 v8
keptrolled backsubmitted
  1. v0The agent inherited a vectorized PyTorch paged-GQA decode baseline1$12.49
  2. v1The agent wrote a fused split-KV Triton flash-decoding kernel27.13159 min · $24.98
  3. v2The agent added a fast-launch plan cache and retuned the combine kernel33.33382 min · $43.24
  4. v3The agent switched split selection to round-to-two-waves and fixed the cache key33.36100 min · $56.30
  5. v4The agent rewrote the combine kernel and added a single-split direct path37.19117 min · $68.54
  6. v5The agent raised the split factor to 2.5 target waves38.946132 min · $77.09
  7. v6The agent floored split chunks at BLOCK_N to stop micro-splits48.015167 min · $101.12
  8. v6 candidateThe agent tried a pre-rotated K workspace and reverted it+4-8% on rope cases$108.05
  9. v6 candidate 2The agent tried f16x2-packed tanh and exp2 PTX opsNaN for cap>=17$114.98
  10. v7The agent quantized chunks to whole tiles and dropped the lse gather51.635186 min · $121.90
  11. v8The agent halved COMBINE_M to eight and folded the denominator inline54.906211 min · $144.16
  12. v8 candidateThe agent tried native-dtype p@V operands and hit fp16 underflowfp16 fails, err ~30-39$146.91

On the hidden set

Original metricNormalised score
Starter1.00x0.00
Frontier-calibrated reference73.55x0.60
Upper233.38x1.00
This run (Qwen3.8 Max)54.08710.4404
236 minWall clock
-Spend
-Tokens
3Versions, 2 kept

On the visible set

1.0 1.5 2.0 2.5 3.0 3.5 0 0 1 2 2 Agent step Visible-panel speedup vs verifie ↑ v0 v1 v2
keptrolled backsubmitted
  1. v0The agent inherited a vectorized PyTorch paged-GQA decode baseline1$0.00
  2. v1The agent added alignment hints so K/V loads vectorized into cp.async1.9202 min · $0.00
  3. v2The agent stored normalized fp16 partials and read KV once per query group3.5$0.00

On the hidden set

Original metricNormalised score
Starter1.00x0.00
Frontier-calibrated reference73.55x0.60
Upper233.38x1.00
This run (GLM 5.3)68.23140.5491
-Wall clock
$11.81Spend
16.7MTokens
20Versions, 8 kept

On the visible set

1.5 3.0 4.5 6.0 0 4 8 12 16 Agent step Visible-panel speedup vs verifie ↑ v0 v0-measured v1 v2 v3 v4 v5 v6 v7 v8 v9 v10 v11 v12 v13 v14 v15 v16 v17 v18
keptrolled backsubmitted
  1. v0The agent inherited a vectorized PyTorch paged-GQA decode baseline1
  2. v0-measuredThe agent measured the inherited baseline on the full visible panel1.0024
  3. v1The agent replaced einsum contractions with torch.matmul on permuted operands1.0018
  4. v2The agent added a Triton streaming softmax kernel for the no-RoPE path1.719
  5. v3The agent added a Triton streaming RoPE path with cached sine/cosine tables2.6996
  6. v4The agent raised the Triton token block size from 32 to 643.2564
  7. v5The agent raised the token block size from 64 to 1282.7955
  8. v6The agent raised the query-head block size from 16 to 322.5225
  9. v7The agent raised Triton num_warps from 4 to 84.1836
  10. v8The agent raised num_warps from 8 to 163.5173
  11. v9The agent cast softmax probabilities to value dtype for p@V4.497
  12. v10The agent cast rotated RoPE Q/K tiles back to input dtype5.7955
  13. v11The agent retested BLOCK_N=128 after the low-precision dot changes5.4526
  14. v12The agent retested BLOCK_N=32 after the low-precision dot changes4.2247
  15. v13The agent stored the cached RoPE tables in the input floating dtype6.7598
  16. v14The agent raised Triton num_stages from 3 to 46.7557
  17. v15The agent lowered Triton num_stages from 3 to 26.7325
  18. v16The agent dropped the no-RoPE kernel to four warps6.5554
  19. v17The agent dropped the RoPE kernel to four warps4.2773
  20. v18The agent added a Triton fast path for window_left equal to zero6.7415

On the hidden set

Original metricNormalised score
Starter1.00x0.00
Frontier-calibrated reference73.55x0.60
Upper233.38x1.00
This run (GPT-5.5)6.66150.0499

Leaderboard

Where each run landed on the sealed held-out set, on the same normalised-score scale as the anchors above.

0 0.3 0.6 1.0 1 Opus 5 claude code · max 0.569 2 GLM 5.3 claude code · max 0.549 3 Kimi K3 kimi cli · max 0.494 4 Qwen3.8 Max qwen coder · xhigh 0.440 5 DeepSeek V4 Pro claude code · max 0.207 6 GPT-5.6-sol codex · max 0.191 7 Gemini 3.7 Flash antigravity · high 0.157 8 Grok 4.6 grok · xhigh 0.125 9 GPT-5.5 codex · xhigh 0.050