diff --git a/docs/report/1H-codesign-paper/build/main.pdf b/docs/report/1H-codesign-paper/build/main.pdf
index 803e246..5e5f73d 100644
Binary files a/docs/report/1H-codesign-paper/build/main.pdf and b/docs/report/1H-codesign-paper/build/main.pdf differ
diff --git a/docs/report/1H-codesign-paper/figures/gqa_short_context/gemm_util_a1_variant_ablation.png b/docs/report/1H-codesign-paper/figures/gqa_short_context/gemm_util_a1_variant_ablation.png
index faf5d94..8c91e8d 100644
Binary files a/docs/report/1H-codesign-paper/figures/gqa_short_context/gemm_util_a1_variant_ablation.png and b/docs/report/1H-codesign-paper/figures/gqa_short_context/gemm_util_a1_variant_ablation.png differ
diff --git a/docs/report/1H-codesign-paper/figures/gqa_short_context/per_cube_tradeoff_mode_compare.png b/docs/report/1H-codesign-paper/figures/gqa_short_context/per_cube_tradeoff_mode_compare.png
index 80a3762..a80f262 100644
Binary files a/docs/report/1H-codesign-paper/figures/gqa_short_context/per_cube_tradeoff_mode_compare.png and b/docs/report/1H-codesign-paper/figures/gqa_short_context/per_cube_tradeoff_mode_compare.png differ
diff --git a/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_a1_variant_compare.png b/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_a1_variant_compare.png
index 3faa027..fc228be 100644
Binary files a/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_a1_variant_compare.png and b/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_a1_variant_compare.png differ
diff --git a/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant1_mode_compare.png b/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant1_mode_compare.png
index 3085394..a391db2 100644
Binary files a/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant1_mode_compare.png and b/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant1_mode_compare.png differ
diff --git a/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant2_mode_compare.png b/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant2_mode_compare.png
index dc4d83c..c6cac7d 100644
Binary files a/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant2_mode_compare.png and b/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant2_mode_compare.png differ
diff --git a/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant3_mode_compare.png b/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant3_mode_compare.png
index ece328c..8dbf7e6 100644
Binary files a/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant3_mode_compare.png and b/docs/report/1H-codesign-paper/figures/gqa_short_context/wall_variant3_mode_compare.png differ
diff --git a/docs/report/1H-codesign-paper/sections/05-gqa.tex b/docs/report/1H-codesign-paper/sections/05-gqa.tex
index 0eda884..6bd8b4e 100644
--- a/docs/report/1H-codesign-paper/sections/05-gqa.tex
+++ b/docs/report/1H-codesign-paper/sections/05-gqa.tex
@@ -141,10 +141,10 @@ CUBE to each head and uses all eight PEs of that CUBE on the head's
 sequence shard; \textsf{2-kv-per-cube} packs two heads per CUBE with
 group size four; \textsf{4-kv-per-cube} packs four heads with group size
 two; \textsf{8-kv-per-cube} loads all heads onto a single CUBE with one
-PE per head. The mapping fixes a trade-off: as KV heads per CUBE grows,
-HBM bandwidth pressure per CUBE grows linearly, the per-CUBE KV cache
-grows linearly, but the intra-CUBE IPCQ broadcast/reduce cost shrinks to
-zero (8-kv-per-cube uses a single PE per head and has no group
+PE per head. The mapping fixes a trade-off: as KV heads per CUBE grows, the per-CUBE KV
+footprint grows linearly and the mapping uses proportionally fewer CUBEs
+(eight down to one), but the intra-CUBE IPCQ broadcast/reduce cost shrinks
+to zero (8-kv-per-cube uses a single PE per head and has no group
 communication).
 
 A second axis is how each GEMM tile is issued. We isolate this with three
@@ -160,101 +160,130 @@ We sweep all four mappings $\times$ three composite tiers $\times$
 $S_{kv}\in\{8\text{K},16\text{K},32\text{K},64\text{K}\}$ for both
 phases (prefill uses $T_q{=}8$ sliced tiles; decode uses $T_q{=}1$). The
 headline latency comparison is in Figure~\ref{fig:gqa-short-wall-baseline}.
-At short context the two mappings are effectively tied, and they separate
-as context grows: the $\textsf{8-kv-per-cube}/\textsf{1-kv-per-cube}$
-latency ratio rises monotonically with $S_{kv}$ (decode: $0.98\times$ at
-$8$K---within measurement noise---to $1.06\times$ at $64$K, where
-\textsf{1-kv-per-cube} is the faster; prefill follows the same trend). The
-wall gap is
-modest because short-context decode is dominated by setup overhead and
-its $M{=}G{=}8$ skinny shape leaves the MAC array idle most of the time;
-where the mappings really separate is per-CUBE pressure.
+The mappings separate cleanly, and in proportion to how many PEs each
+dedicates to a head: \textsf{1-kv-per-cube} spreads one head across all
+eight PEs of its CUBE, \textsf{8-kv-per-cube} gives each head a single PE,
+and the four mappings form a $\{64,32,16,8\}$-active-PE ladder. Decode
+latency tracks that ladder inversely---at $S_{kv}{=}64$K the
+\textsf{8-kv-per-cube}/\textsf{1-kv-per-cube} ratio is $7.4\times$
+($\SI{87.0}{}$ vs.\ $\SI{11.8}{\micro\second}$), and even at $8$K it is
+$4.8\times$ ($\SI{10.8}{}$ vs.\ $\SI{2.2}{\micro\second}$); prefill shows
+the same ordering more gently ($2.3\times$ at $64$K). The reason is that
+decode here is \emph{bandwidth-bound}: each active PE streams its KV shard
+at \SI{46}{}--\SI{76}{\percent} of the \SI{256}{\giga\byte\per\second}
+per-PE ceiling (rising with context), so a mapping's speed is set by how
+many PEs it puts to work---more PEs, more aggregate HBM bandwidth, lower
+latency. The $M{=}G{=}8$ score GEMM is skinny and leaves the MAC array
+lightly loaded throughout; the lever here is data movement, not compute.
 
 \begin{figure}[t]
 \centering
 \includegraphics[width=\linewidth]{gqa_short_context/wall_variant1_mode_compare.png}
 \caption{Wall-clock latency of the four short-context mappings, baseline
-kernel (no composite). Prefill and decode follow the same ordering; the
-two mappings are effectively tied at $8$K and separate as context grows,
-the \textsf{8-kv-per-cube}/\textsf{1-kv-per-cube} ratio widening from
-$0.98\times$ at $8$K to $1.06\times$ at $64$K (\textsf{1-kv-per-cube}
-faster at the large end). The composite and
-softmax\_merge tiers track the baseline within a few percent at this
-shape (see Figure~\ref{fig:gqa-short-a1-variant}).}
+kernel (no composite; log scale). The mappings separate along the
+$\{64,32,16,8\}$-active-PE ladder set by heads-per-CUBE:
+\textsf{1-kv-per-cube} (64 PEs) is fastest and \textsf{8-kv-per-cube}
+(8 PEs) slowest, by $4.8\times$ at $8$K growing to $7.4\times$ at $64$K in
+decode ($2.3\times$ at $64$K in prefill). Decode is bandwidth-bound, so
+latency scales inversely with the number of active PEs. Composite tiers
+track the baseline at this skinny shape
+(Figure~\ref{fig:gqa-short-a1-variant}).}
 \label{fig:gqa-short-wall-baseline}
 \end{figure}
 
-Figure~\ref{fig:gqa-short-tradeoff} makes the per-CUBE trade-off
-explicit. Per-PE HBM bandwidth utilization is $\sim$\,$8\times$ higher on
-\textsf{8-kv-per-cube} than on \textsf{1-kv-per-cube} at every context
-length: putting all heads on a single CUBE saturates that CUBE's HBM
-channel while leaving seven CUBEs idle, whereas spreading one head per
-CUBE keeps each CUBE's read bandwidth well below the
-\SI{256}{\giga\byte\per\second} per-PE ceiling. IPCQ traffic moves in the
-opposite direction---\textsf{1-kv-per-cube} pays the broadcast or
-chain-reduce cost across all eight PEs, while \textsf{8-kv-per-cube} pays
-nothing---but absolute IPCQ volume stays under \SI{120}{\kibi\byte} per
-run, two orders of magnitude below the HBM traffic. Because per-PE HBM
-bandwidth and per-CUBE KV cache (which grows as $\text{kv\_per\_cube}\times
-S_{kv}\times d_{\text{head}}$, so $8\times$ more on
-\textsf{8-kv-per-cube}) both scale with the number of heads on a CUBE,
-the small wall advantage of \textsf{1-kv-per-cube} compounds with much
-larger headroom for longer contexts and for batched serving---which is
-why it is the choice carried into long context (\S\ref{sec:gqa-long}).
+Figure~\ref{fig:gqa-short-tradeoff} shows what the mappings trade for that
+latency. The differentiator is \emph{density}: \textsf{8-kv-per-cube}
+concentrates all eight heads onto one CUBE, so its per-CUBE KV footprint is
+$8\times$ that of \textsf{1-kv-per-cube} (one CUBE per head), and it
+occupies a single one of the SIP's sixteen CUBEs against
+\textsf{1-kv-per-cube}'s eight. Per-PE HBM \emph{utilization}, by contrast,
+is similar across mappings---every active PE runs near the same
+\SI{46}{}--\SI{76}{\percent} of the per-PE ceiling---so the latency ladder
+comes from the \emph{count} of active PEs, not from any per-PE
+concentration. IPCQ traffic runs opposite to density:
+\textsf{1-kv-per-cube} pays the eight-PE chain-reduce while
+\textsf{8-kv-per-cube} (one PE per head) pays none, but absolute IPCQ volume
+stays under \SI{120}{\kibi\byte} per run, two orders of magnitude below the
+HBM traffic. The result is a genuine trade-off rather than a single winner:
+\textsf{1-kv-per-cube} minimizes per-request latency by spending the most
+hardware (eight CUBEs, 64 PEs) on one request, while denser mappings leave
+CUBEs free for other requests---the batched-serving axis taken up below.
 
 \begin{figure}[t]
 \centering
 \includegraphics[width=\linewidth]{gqa_short_context/per_cube_tradeoff_mode_compare.png}
 \caption{Per-CUBE trade-off across the four mappings (baseline kernel;
-HBM and IPCQ totals are variant-invariant). Top row: per-PE HBM
-bandwidth utilization---\textsf{8-kv-per-cube} concentrates roughly
-$8\times$ more bandwidth on a single CUBE than \textsf{1-kv-per-cube}.
-Bottom row: IPCQ traffic per run on a log scale---\textsf{1-kv-per-cube}
-pays the broadcast/reduce cost across the largest group while
-\textsf{8-kv-per-cube} (single PE per head) pays zero.}
+log scale). Top row: per-CUBE KV footprint---\textsf{8-kv-per-cube} packs
+all eight heads onto one CUBE, an $8\times$ larger per-CUBE KV than
+\textsf{1-kv-per-cube} and one CUBE used against eight (the $8\times$ ratio
+is context-invariant; the absolute footprint grows with $S_{kv}$). Bottom
+row: IPCQ traffic per run---\textsf{1-kv-per-cube} pays the eight-PE
+chain-reduce while \textsf{8-kv-per-cube} (single PE per head) pays zero.
+The density (top) is what sets how many concurrent requests a SIP can hold.}
 \label{fig:gqa-short-tradeoff}
 \end{figure}
 
-The composite ablation in Figure~\ref{fig:gqa-short-a1-variant} shows
-the three tiers within \textsf{1-kv-per-cube}: wall-clock is essentially
-unchanged across tiers ($< 0.5\%$ in either direction). This is the
-expected outcome for the decode-skinny shape---$M{=}G{=}8$ is well below
-the scheduler's $\textsf{TILE\_M}{=}32$ supertile (\S\ref{sec:gemm}), so
-the composite path pads $M$ by $4\times$ with zeros and the fusion has
-no slack to win back. The padding shows up in
-Figure~\ref{fig:gqa-short-gemm-util} as inflated GEMM engine utilization
-on the composite and softmax\_merge tiers (roughly $3$ and $6\times$ the
-baseline), which is bookkeeping cost rather than added useful work.
-Demonstrating fusion benefit therefore requires lifting $M$---either by
-batching multiple users (a separate batched-$M$ kernel, deferred to
-\S\ref{sec:future}) or by the long-context path, where Q-tile splitting
-yields a larger $M$ per PE. In the short-context regime the contribution
-of composite/softmax\_merge is a no-regression dispatch-cost isolation;
-the substantive design lever is the mapping, and
-\textsf{1-kv-per-cube} wins it.
+The composite ablation in Figure~\ref{fig:gqa-short-a1-variant} shows the
+three tiers within \textsf{1-kv-per-cube}: wall-clock is nearly flat, with
+the GEMM-only composite tier a few percent \emph{slower}
+(\SI{14.0}{} vs.\ \SI{11.8}{\micro\second} at $64$K decode) and the
+\textsf{softmax\_merge} tier matching the baseline. This is the expected
+outcome for the decode-skinny shape---$M{=}G{=}8$ is well below the
+scheduler's $\textsf{TILE\_M}{=}32$ supertile (\S\ref{sec:gemm}), so the
+composite path pads $M$ by $4\times$ with zeros and the fusion has no slack
+to win back; the padding is pure overhead here. It shows up in
+Figure~\ref{fig:gqa-short-gemm-util} as inflated GEMM-engine busy time on
+the composite tiers---the padded-$M$ GEMM is charged in full against a
+sub-\SI{15}{\micro\second} decode wall, driving the accounted utilization
+well above the baseline's \SI{12}{}--\SI{18}{\percent}. Fusion benefit
+requires lifting $M$, either by the long-context Q-tile split
+(\S\ref{sec:gqa-composite}) or by batching multiple users---which is also
+the lever that converts a mapping's density into throughput.
+
+\paragraph{Latency versus batched throughput (projected).} The latency
+ladder and the density trade-off pull in opposite directions once a SIP
+serves more than one request. Decode users are independent---each attends
+its own KV cache---and these mappings generate no cross-CUBE traffic, so a
+SIP can run several users at once on disjoint CUBE groups, up to $16/C$
+users: two for \textsf{1-kv-per-cube}, sixteen for \textsf{8-kv-per-cube}.
+Projecting from the measured single-user latencies under perfect overlap
+(an idealization---no concurrent multi-user run has been measured yet; see
+\S\ref{sec:future}), aggregate throughput at a full SIP is
+$\text{users}/\text{latency}$: at $8$K it rises from \SI{0.89}{}
+(\textsf{1-kv-per-cube}, 2 users) to \SI{1.48}{}~requests\,/\,\si{\micro\second}
+(\textsf{8-kv-per-cube}, 16 users), so the \emph{dense} mapping wins on
+throughput even though it is the \emph{slowest} per request; at $64$K the
+projected throughputs converge (\SI{0.17}{}--\SI{0.18}{}) because latency
+there already scales as $1/C$. The design conclusion is therefore
+regime-dependent: \textsf{1-kv-per-cube} is the latency-optimal choice for
+single-stream or low-batch decode, while denser mappings are the
+throughput-optimal choice for high-batch short-context serving. These are
+\emph{projections} assuming zero cross-user interference; validating them
+against a concurrent multi-user run is 2H work (\S\ref{sec:future}).
 
 \begin{figure}[t]
 \centering
 \includegraphics[width=\linewidth]{gqa_short_context/wall_a1_variant_compare.png}
-\caption{Composite-tier ablation on \textsf{1-kv-per-cube}: wall-clock
-latency is flat across the three tiers within $0.5\%$. With
-$M{=}G{=}8 < \textsf{TILE\_M}{=}32$, the composite scheduler pads the
-$M$ dimension $4\times$ and any fusion overlap is absorbed by the
-padding overhead. Fusion benefit is exercised only when $M$ fills the
-supertile---batched-$M$ inference or long-context Q-tile splitting.}
+\caption{Composite-tier ablation on \textsf{1-kv-per-cube} (log scale):
+wall-clock is nearly flat across tiers---the GEMM-only composite runs a few
+percent slower (padding overhead) and \textsf{softmax\_merge} matches the
+baseline. With $M{=}G{=}8 < \textsf{TILE\_M}{=}32$ the composite scheduler
+pads $M$ by $4\times$, leaving no fusion slack to recover at this shape.
+Fusion pays off only when $M$ fills the supertile---long-context Q-tile
+splitting or batched-$M$ inference.}
 \label{fig:gqa-short-a1-variant}
 \end{figure}
 
 \begin{figure}[t]
 \centering
 \includegraphics[width=\linewidth]{gqa_short_context/gemm_util_a1_variant_ablation.png}
-\caption{GEMM engine utilization on \textsf{1-kv-per-cube} across the
-three composite tiers. The $\sim 3$ and $\sim 6\times$ rise on the
-composite and softmax\_merge tiers reflects supertile padding (the
-$8\!\to\!32$ pad in $M$) being accounted as GEMM time, not added useful
-work. With actual $M=8$, the engine is below \SI{2}{\percent} busy in
-all tiers---the kernel is firmly setup- and KV-bandwidth-bound at this
-shape.}
+\caption{GEMM-engine busy fraction on \textsf{1-kv-per-cube} across the
+three composite tiers. The baseline runs at \SI{12}{}--\SI{18}{\percent};
+the composite tiers read much higher---up to and past
+\SI{100}{\percent} of the short decode wall---because the padded
+$8\!\to\!32$ $M$ dimension is charged as GEMM time, bookkeeping overhead
+rather than useful work. The engine is never the bottleneck at this shape;
+decode is KV-bandwidth-bound.}
 \label{fig:gqa-short-gemm-util}
 \end{figure}
 
diff --git a/docs/sweeps/plot_scripts/plot_short_context.py b/docs/sweeps/plot_scripts/plot_short_context.py
index 6ae61bb..728447a 100644
--- a/docs/sweeps/plot_scripts/plot_short_context.py
+++ b/docs/sweeps/plot_scripts/plot_short_context.py
@@ -23,7 +23,7 @@ from pathlib import Path
 import matplotlib.pyplot as plt
 import numpy as np
 
-ROOT = Path("/Users/lalicorne/Desktop/ahbm/kernbench2")
+ROOT = Path(__file__).resolve().parents[3]
 SWEEPS = ROOT / "docs" / "sweeps"
 OUT = (ROOT / "docs" / "report" / "1H-codesign-paper"
        / "figures" / "gqa_short_context")
@@ -156,8 +156,8 @@ def fig_wall_a1_variant_compare(fname):
 def fig_per_cube_tradeoff(fname):
     fig, axes = plt.subplots(2, 2, figsize=(13, 9))
     metrics = [
-        ("hbm_bw_util", "HBM BW utilization (per PE)", False, "{:.3f}"),
-        ("ipcq_kb",     "IPCQ traffic (KB)",           True,  "{:.1f}"),
+        ("kv_cache_per_cube_mb", "Per-cube KV footprint (MB)", True, "{:.0f}"),
+        ("ipcq_kb",              "IPCQ traffic (KB)",          True, "{:.1f}"),
     ]
     for row, (key, ylabel, log, fmt) in enumerate(metrics):
         for col, phase in enumerate(("prefill", "decode")):
diff --git a/docs/sweeps/short_context_decode_composite_fused_sweep.csv b/docs/sweeps/short_context_decode_composite_fused_sweep.csv
index 19294ce..d62c2c1 100644
--- a/docs/sweeps/short_context_decode_composite_fused_sweep.csv
+++ b/docs/sweeps/short_context_decode_composite_fused_sweep.csv
@@ -1,17 +1,17 @@
 variant,mode,kv_per_cube,C,S_kv,wall_us,n_pe,gemm_util,math_count,math_pipeline_us,hbm_bw_util,hbm_read_mb,hbm_write_kb,ipcq_kb,kv_cache_per_cube_mb
-composite_fused,A1,1,8,8192,146.95153450000072,64,0.0017838806576053131,656,0.0,0.013997812319543966,16.0625,8.0,57.75,2.0
-composite_fused,A1,1,8,16384,279.867342500003,64,0.007173570099630665,1296,13.903615999992937,0.0146676634841736,32.0625,8.0,57.75,4.0
-composite_fused,A1,1,8,32768,544.5763825000042,64,0.010097125355742134,2576,41.71084799996763,0.015059411798858053,64.0625,8.0,57.75,8.0
-composite_fused,A1,1,8,65536,1074.101862500392,64,0.011619635377527878,5136,97.32531200143696,0.015262053416273654,128.0625,8.0,57.75,16.0
-composite_fused,A2,2,4,8192,140.35490875000028,32,0.01430408111743636,624,6.951807999995537,0.029254409671653128,16.03125,8.0,24.75,4.0
-composite_fused,A2,2,4,16384,273.9919487500021,32,0.020068677291975947,1264,20.855423999989405,0.029935186188568382,32.03125,8.0,24.75,8.0
-composite_fused,A2,2,4,32768,541.3734287504049,32,0.0230537210304644,2544,48.66265599996224,0.030282239817052973,64.03125,8.0,24.75,16.0
-composite_fused,A2,2,4,65536,1076.851308751879,32,0.024557433125924955,5104,104.27712000153959,0.030438742780552714,128.03125,8.0,24.75,32.0
-composite_fused,A4,4,2,8192,140.2291675000001,16,0.03921192786066565,608,10.427711999993306,0.058504233792873325,16.015625,8.0,8.25,8.0
-composite_fused,A4,4,2,16384,276.4849475001752,16,0.04514051167338893,1248,24.33132799998764,0.05930160085835996,32.015625,8.0,8.25,16.0
-composite_fused,A4,4,2,32768,549.3539675010354,16,0.04813782290440163,2528,52.13855999995954,0.059670088757369746,64.015625,8.0,8.25,32.0
-composite_fused,A4,4,2,65536,1095.4730075019588,16,0.04963405545386165,5088,107.7530240015909,0.05983534012350623,128.015625,8.0,8.25,64.0
-composite_fused,B,8,1,8192,144.10828500005923,8,0.08660620726926088,600,12.16566399999219,0.11380331116974475,16.0078125,8.0,0.0,16.0
-composite_fused,B,8,1,16384,285.6728750004253,8,0.09256988084807763,1240,26.069279999986758,0.11476063311909189,32.0078125,8.0,0.0,32.0
-composite_fused,B,8,1,32768,568.9925550010347,8,0.09555971782205692,2520,53.8765119999582,0.1152071313127828,64.0078125,8.0,0.0,64.0
-composite_fused,B,8,1,65536,1135.6319150019874,8,0.09706392938508733,5080,109.49097600161657,0.11543176822374755,128.0078125,8.0,0.0,128.0
+composite_fused,A1,1,8,8192,2.2471645000005376,64,0.11665545624270428,656,0.0,0.457688,16.0625,8.0,57.75,2.0
+composite_fused,A1,1,8,16384,4.090972500001895,64,0.4907507933676371,1296,13.903615999992937,0.501714,32.0625,8.0,57.75,4.0
+composite_fused,A1,1,8,32768,6.65601250000007,64,0.8261186408347434,2576,41.71084799996763,0.61606,64.0625,8.0,57.75,8.0
+composite_fused,A1,1,8,65536,11.786092499984429,64,1.05893212704674,5136,97.32531200143696,0.695438,128.0625,8.0,57.75,16.0
+composite_fused,A2,2,4,8192,3.8395887500001407,32,0.5228809986460942,624,6.951807999995537,0.534693,16.03125,8.0,24.75,4.0
+composite_fused,A2,2,4,16384,6.404628750001605,32,0.858544064722257,1264,20.855423999989405,0.640318,32.03125,8.0,24.75,8.0
+composite_fused,A2,2,4,32768,11.534708749999874,32,1.0820101547616419,2544,48.66265599996224,0.710638,64.03125,8.0,24.75,16.0
+composite_fused,A2,2,4,65536,21.794868749984772,32,1.21334541192213,5104,104.27712000153959,0.751966,128.03125,8.0,24.75,32.0
+composite_fused,A4,4,2,8192,5.915787499999919,16,0.9294884239792724,608,10.427711999993306,0.693399,16.015625,8.0,8.25,8.0
+composite_fused,A4,4,2,16384,11.045867500001041,16,1.1298951395303998,1248,24.33132799998764,0.742178,32.015625,8.0,8.25,16.0
+composite_fused,A4,4,2,32768,21.306027499999384,16,1.2411841672219757,2528,52.13855999995954,0.769266,64.015625,8.0,8.25,32.0
+composite_fused,A4,4,2,65536,41.82634749998455,16,1.2999645260103314,5088,107.7530240015909,0.783573,128.015625,8.0,8.25,64.0
+composite_fused,B,8,1,8192,10.856904999999854,8,1.149560763397397,600,12.16566399999219,0.75528,16.0078125,8.0,0.0,16.0
+composite_fused,B,8,1,16384,21.117065000001226,8,1.2522906947682175,1240,26.069279999986758,0.776244,32.0078125,8.0,0.0,32.0
+composite_fused,B,8,1,32768,41.637384999999774,8,1.3058641410537761,2520,53.8765119999582,0.787177,64.0078125,8.0,0.0,64.0
+composite_fused,B,8,1,65536,82.67802499997313,8,1.33323087973183,5080,109.49097600161657,0.792762,128.0078125,8.0,0.0,128.0
diff --git a/docs/sweeps/short_context_decode_composite_sweep.csv b/docs/sweeps/short_context_decode_composite_sweep.csv
index 2c70e9d..437f754 100644
--- a/docs/sweeps/short_context_decode_composite_sweep.csv
+++ b/docs/sweeps/short_context_decode_composite_sweep.csv
@@ -1,17 +1,17 @@
 variant,mode,kv_per_cube,C,S_kv,wall_us,n_pe,gemm_util,math_count,math_pipeline_us,hbm_bw_util,hbm_read_mb,hbm_write_kb,ipcq_kb,kv_cache_per_cube_mb
-composite,A1,1,8,8192,147.23658450000076,64,0.0057446592018384314,656,0.0,0.013970712557516501,16.0625,8.0,57.75,2.0
-composite,A1,1,8,16384,279.9533665000026,64,0.006042606385378475,1360,0.0,0.014663156408229735,32.0625,8.0,57.75,4.0
-composite,A1,1,8,32768,545.3869305000062,64,0.006203478321011016,2768,0.0,0.015037030668265908,64.0625,8.0,57.75,8.0
-composite,A1,1,8,65536,1076.3614585003581,64,0.006286542449928883,5584,0.0,0.015230013923798023,128.0625,8.0,57.75,16.0
-composite,A2,2,4,8192,140.44093275000077,32,0.012045263206857853,656,0.0,0.02923649052736712,16.03125,8.0,24.75,4.0
-composite,A2,2,4,16384,274.80249675000204,32,0.012311736756599198,1360,0.0,0.029846890392199246,32.03125,8.0,24.75,8.0
-composite,A2,2,4,32768,543.6330247504077,32,0.012446984807344498,2768,0.0,0.030156372504276757,64.03125,8.0,24.75,16.0
-composite,A2,2,4,65536,1082.0090007518106,32,0.012507459726542472,5584,0.0,0.030293648183356066,128.03125,8.0,24.75,32.0
-composite,A4,4,2,8192,141.0397155000009,16,0.023988250316491475,656,0.0,0.05816801296653173,16.015625,8.0,8.25,8.0
-composite,A4,4,2,16384,278.7445435001758,16,0.02427524469220757,1360,0.0,0.05882088235384476,32.015625,8.0,8.25,16.0
-composite,A4,4,2,32768,554.51165950104,16,0.024405589617097562,2768,0.0,0.059115077994024615,64.015625,8.0,8.25,32.0
-composite,A4,4,2,65536,1106.4268915018213,16,0.024462861675913012,5584,0.0,0.05924295631591859,128.015625,8.0,8.25,64.0
-composite,B,8,1,8192,146.3478710000607,8,0.0462363541999099,656,0.0,0.11206176002378058,16.0078125,8.0,0.0,16.0
-composite,B,8,1,16384,290.8105570004271,8,0.046536082251474285,1360,0.0,0.1127331838917796,32.0078125,8.0,0.0,32.0
-composite,B,8,1,32768,579.9264290010428,8,0.046672071914328234,2768,0.0,0.11303502775846438,64.0078125,8.0,0.0,64.0
-composite,B,8,1,65536,1158.158173001712,8,0.04674036523209764,5584,0.0,0.11318661220534877,128.0078125,8.0,0.0,128.0
+composite,A1,1,8,8192,2.532214500000584,64,0.33402541530150626,656,0.0,0.406166,16.0625,8.0,57.75,2.0
+composite,A1,1,8,16384,4.176996500001463,64,0.40499148132408613,1360,0.0,0.491382,32.0625,8.0,57.75,4.0
+composite,A1,1,8,32768,7.466560500002,64,0.4531264428807111,2768,0.0,0.549182,64.0625,8.0,57.75,8.0
+composite,A1,1,8,65536,14.045688499950572,64,0.4817558071541797,5584,0.0,0.58356,128.0625,8.0,57.75,16.0
+composite,A2,2,4,8192,3.9256127500006404,32,0.4309258471789165,656,0.0,0.522976,16.03125,8.0,24.75,4.0
+composite,A2,2,4,16384,7.215176750001498,32,0.4689138072801237,1360,0.0,0.568385,32.03125,8.0,24.75,8.0
+composite,A2,2,4,32768,13.794304750002688,32,0.4905351971318266,2768,0.0,0.594231,64.03125,8.0,24.75,16.0
+composite,A2,2,4,65536,26.95256074991636,32,0.5021112511805295,5584,0.0,0.608068,128.03125,8.0,24.75,32.0
+composite,A4,4,2,8192,6.726335500000743,16,0.5029924540606681,656,0.0,0.609842,16.015625,8.0,8.25,8.0
+composite,A4,4,2,16384,13.305463500001585,16,0.5085574057667106,1360,0.0,0.616138,32.015625,8.0,8.25,16.0
+composite,A4,4,2,32768,26.463719500003965,16,0.5113863151276257,2768,0.0,0.619338,64.015625,8.0,8.25,32.0
+composite,A4,4,2,65536,52.78023149984703,16,0.5128126048745716,5584,0.0,0.620952,128.015625,8.0,8.25,64.0
+composite,B,8,1,8192,13.096491000001318,8,0.5166721375947838,656,0.0,0.626122,16.0078125,8.0,0.0,16.0
+composite,B,8,1,16384,26.254747000003057,8,0.5154566524738308,1360,0.0,0.624344,32.0078125,8.0,0.0,32.0
+composite,B,8,1,32768,52.57125900000788,8,0.5148510519664782,2768,0.0,0.623459,64.0078125,8.0,0.0,64.0
+composite,B,8,1,65536,105.20428299969761,8,0.514548785079354,5584,0.0,0.623016,128.0078125,8.0,0.0,128.0
diff --git a/docs/sweeps/short_context_decode_sweep.csv b/docs/sweeps/short_context_decode_sweep.csv
index 89bfff5..17ac51e 100644
--- a/docs/sweeps/short_context_decode_sweep.csv
+++ b/docs/sweeps/short_context_decode_sweep.csv
@@ -1,17 +1,17 @@
 mode,kv_per_cube,C,S_kv,wall_us,n_pe,gemm_util,math_count,hbm_bw_util,hbm_read_mb,hbm_write_kb,ipcq_kb,kv_cache_per_cube_mb
-A1,1,8,8192,146.95153450000072,64,0.0017838806576053131,656,0.013997812319543966,16.0625,8.0,57.75,2.0
-A1,1,8,16384,279.3832665000025,64,0.0018765905580817512,1360,0.014693077546933053,32.0625,8.0,57.75,4.0
-A1,1,8,32768,544.246730500006,64,0.0019266555796063503,2768,0.015068533333155061,64.0625,8.0,57.75,8.0
-A1,1,8,65536,1074.081058500367,64,0.0019525081309290412,5584,0.01526234902874828,128.0625,8.0,57.75,16.0
-A2,2,4,8192,139.87083275000066,32,0.003748372621310069,656,0.029355655637933423,16.03125,8.0,24.75,4.0
-A2,2,4,16384,273.6622967500018,32,0.0038316421825465713,1360,0.029971245938539923,32.03125,8.0,24.75,8.0
-A2,2,4,32768,541.3526247504072,32,0.0038739112070761647,2768,0.030283403553383558,64.03125,8.0,24.75,16.0
-A2,2,4,65536,1077.4482007518286,32,0.003892812663356761,5584,0.030421880121130614,128.03125,8.0,24.75,32.0
-A4,4,2,8192,139.89951550000075,16,0.007495208230366452,656,0.05864209015076937,16.015625,8.0,8.25,8.0
-A4,4,2,16384,276.4641435001754,16,0.007585620230706115,1360,0.05930606331952627,32.015625,8.0,8.25,16.0
-A4,4,2,32768,549.9508595010393,16,0.007626688689616412,2768,0.059605325518975856,64.015625,8.0,8.25,32.0
-A4,4,2,65536,1097.305291501857,16,0.0076447348472311214,5584,0.059735426874945555,128.015625,8.0,8.25,64.0
-B,8,1,8192,144.06747100006032,8,0.014556735017573057,656,0.11383555139933797,16.0078125,8.0,0.0,16.0
-B,8,1,16384,286.24975700042637,8,0.014652602831705778,1360,0.1145293548666704,32.0078125,8.0,0.0,32.0
-B,8,1,32768,570.8048290010414,8,0.014696105522939606,2768,0.11484135499470417,64.0078125,8.0,0.0,64.0
-B,8,1,65536,1139.9149730017834,8,0.014717953880200899,5584,0.11499805082373886,128.0078125,8.0,0.0,128.0
+A1,1,8,8192,2.2471645000005376,64,0.11665545624270428,656,0.457688,16.0625,8.0,57.75,2.0
+A1,1,8,16384,3.6068965000013704,64,0.14535709577464812,1360,0.569049,32.0625,8.0,57.75,4.0
+A1,1,8,32768,6.326360500001814,64,0.1657471147905734,2768,0.648161,64.0625,8.0,57.75,8.0
+A1,1,8,65536,11.765288499959512,64,0.17824909265982342,5584,0.696668,128.0625,8.0,57.75,16.0
+A2,2,4,8192,3.3555127500005475,32,0.15624676139283236,656,0.611829,16.03125,8.0,24.75,4.0
+A2,2,4,16384,6.074976750001311,32,0.17260576347056114,1360,0.675064,32.03125,8.0,24.75,8.0
+A2,2,4,32768,11.513904750002315,32,0.18214081543451005,2768,0.711922,64.03125,8.0,24.75,16.0
+A2,2,4,65536,22.391760749934242,32,0.18731461303283142,5584,0.731921,128.03125,8.0,24.75,32.0
+A4,4,2,8192,5.586135500000557,16,0.1877104484844272,656,0.734318,16.015625,8.0,8.25,8.0
+A4,4,2,16384,11.025063500001211,16,0.19021677290108474,1360,0.743578,32.015625,8.0,8.25,16.0
+A4,4,2,32768,21.902919500003218,16,0.19149520227204342,2768,0.748302,64.015625,8.0,8.25,32.0
+A4,4,2,65536,43.65863149988279,16,0.19214088284049563,5584,0.750688,128.015625,8.0,8.25,64.0
+B,8,1,8192,10.816091000000947,8,0.1938918598225168,656,0.75813,16.0078125,8.0,0.0,16.0
+B,8,1,16384,21.693947000002314,8,0.1933398288471476,1360,0.755602,32.0078125,8.0,0.0,32.0
+B,8,1,32768,43.449659000006385,8,0.19306499045255035,2768,0.754344,64.0078125,8.0,0.0,64.0
+B,8,1,65536,86.96108299976913,8,0.19292786406575965,5584,0.753716,128.0078125,8.0,0.0,128.0
diff --git a/docs/sweeps/short_context_prefill_composite_fused_sweep.csv b/docs/sweeps/short_context_prefill_composite_fused_sweep.csv
index 2f365bc..192d520 100644
--- a/docs/sweeps/short_context_prefill_composite_fused_sweep.csv
+++ b/docs/sweeps/short_context_prefill_composite_fused_sweep.csv
@@ -1,17 +1,17 @@
 variant,mode,kv_per_cube,C,S_kv,wall_us,n_pe,gemm_util,math_count,math_pipeline_us,hbm_bw_util,hbm_read_mb,hbm_write_kb,ipcq_kb,kv_cache_per_cube_mb
-composite_fused,A1,1,8,8192,1101.4436329998505,64,0.011331194467560951,4800,97.32531200143696,0.0018739043362378583,16.0625,64.0,114688.0,2.0
-composite_fused,A1,1,8,16384,2178.44091796638,64,0.012139279878164511,9920,208.55424000307917,0.0018875884886695197,32.0625,64.0,229376.0,4.0
-composite_fused,A1,1,8,32768,4335.060767847123,64,0.012542561892233649,20160,431.01209600543973,0.0018933990639481287,64.0625,64.0,458752.0,8.0
-composite_fused,A1,1,8,65536,8649.804197613514,64,0.012743513431742019,40640,875.9278079059123,0.0018959966752224023,128.0625,64.0,917504.0,16.0
-composite_fused,A2,2,4,8192,1093.9720835019473,32,0.01853321150177446,2400,96.40870400077105,0.0037734052470386023,16.0625,64.0,49152.0,4.0
-composite_fused,A2,2,4,16384,2175.4760434340433,32,0.019695389489835093,4960,206.59008000165224,0.00378032202414797,32.0625,64.0,98304.0,8.0
-composite_fused,A2,2,4,32768,4338.482933349601,32,0.020281493171481216,10080,426.9528320015669,0.003783811127574436,64.0625,64.0,196608.0,16.0
-composite_fused,A2,2,4,65536,8664.499802991075,32,0.020575802877008553,20320,867.678335950613,0.0037855618611332994,128.0625,64.0,393216.0,32.0
-composite_fused,A4,4,2,8192,1102.9720360028073,16,0.03256574312686088,1200,92.28038400042057,0.007485230568419403,16.0625,64.0,16384.0,8.0
-composite_fused,A4,4,2,16384,2196.357595426575,16,0.03449854256839416,2480,197.74368000090124,0.007488762319145705,32.0625,64.0,32768.0,16.0
-composite_fused,A4,4,2,32768,4383.130075368471,16,0.03547127219906801,5040,408.6702720000148,0.00749053745507198,64.0625,64.0,65536.0,32.0
-composite_fused,A4,4,2,65536,8756.675034884938,16,0.03595923415220853,10160,830.5234559737444,0.007491427937962982,128.0625,64.0,131072.0,64.0
-composite_fused,B,8,1,8192,1125.1019960018205,8,0.12366810519999533,712,121.6138560003899,0.014676002761240575,16.0625,64.0,0.0,16.0
-composite_fused,B,8,1,16384,2241.8354359431773,8,0.13192705729555512,1480,260.6011200008355,0.0146736907948643,32.0625,64.0,0.0,32.0
-composite_fused,B,8,1,32768,4475.302315825972,8,0.1360796989728368,3016,538.575647996068,0.014672528326811126,64.0625,64.0,0.0,64.0
-composite_fused,B,8,1,65536,8942.236075603707,8,0.13816184916758104,6088,1094.5247039788662,0.014671945460927953,128.0625,64.0,0.0,128.0
+composite_fused,A1,1,8,8192,35.53158299944201,64,0.3512557265116541,4800,97.32531200143696,0.0290446,16.0625,64.0,114688.0,2.0
+composite_fused,A1,1,8,16384,63.02314799463935,64,0.4196030322618738,9920,208.55424000307917,0.0326229,32.0625,64.0,229376.0,4.0
+composite_fused,A1,1,8,32768,118.00299802135117,64,0.4607744625054221,20160,431.01209600543973,0.0347788,64.0625,64.0,458752.0,8.0
+composite_fused,A1,1,8,65536,227.95666800683736,64,0.4835519703723695,40640,875.9278079059123,0.0359717,128.0625,64.0,917504.0,16.0
+composite_fused,A2,2,4,8192,35.27508350025257,32,0.5747630902258046,2400,96.40870400077105,0.0585116,16.0625,64.0,49152.0,4.0
+composite_fused,A2,2,4,16384,65.58160349091561,32,0.653336388872743,4960,206.59008000165224,0.0627005,32.0625,64.0,98304.0,8.0
+composite_fused,A2,2,4,32768,126.19361352364719,32,0.69726913692688,10080,426.9528320015669,0.0650429,64.0625,64.0,196608.0,16.0
+composite_fused,A2,2,4,65536,247.42072338639758,32,0.7205501525262498,20320,867.678335950613,0.0662839,128.0625,64.0,393216.0,32.0
+composite_fused,A4,4,2,8192,45.62825600103196,16,0.7872118539828873,1200,92.28038400042057,0.0904703,16.0625,64.0,16384.0,8.0
+composite_fused,A4,4,2,16384,87.81637548340065,16,0.8628360665552867,2480,197.74368000090124,0.09365,32.0625,64.0,32768.0,16.0
+composite_fused,A4,4,2,32768,172.19397554247547,16,0.9029073142513345,5040,408.6702720000148,0.0953343,64.0625,64.0,65536.0,32.0
+composite_fused,A4,4,2,65536,340.9491752808783,16,0.9235491703853554,10160,830.5234559737444,0.096202,128.0625,64.0,131072.0,64.0
+composite_fused,B,8,1,8192,68.33786600000411,8,2.036048828365075,712,121.6138560003899,0.120811,16.0625,64.0,0.0,16.0
+composite_fused,B,8,1,16384,133.87386600000133,8,2.209234414765302,1480,260.6011200008355,0.122862,32.0625,64.0,0.0,32.0
+composite_fused,B,8,1,32768,264.9458659999771,8,2.2985744263323977,3016,538.575647996068,0.12392,64.0625,64.0,0.0,64.0
+composite_fused,B,8,1,65536,527.0898659999762,8,2.3439567929361727,6088,1094.5247039788662,0.124457,128.0625,64.0,0.0,128.0
diff --git a/docs/sweeps/short_context_prefill_composite_sweep.csv b/docs/sweeps/short_context_prefill_composite_sweep.csv
index 9192d2a..5ec3d03 100644
--- a/docs/sweeps/short_context_prefill_composite_sweep.csv
+++ b/docs/sweeps/short_context_prefill_composite_sweep.csv
@@ -1,17 +1,17 @@
 variant,mode,kv_per_cube,C,S_kv,wall_us,n_pe,gemm_util,math_count,math_pipeline_us,hbm_bw_util,hbm_read_mb,hbm_write_kb,ipcq_kb,kv_cache_per_cube_mb
-composite,A1,1,8,8192,1101.224596999844,64,0.006144606666764116,5248,0.0,0.0018742770599413814,16.0625,64.0,114688.0,2.0
-composite,A1,1,8,16384,2178.2218819663735,64,0.006212950164859459,10880,0.0,0.0018877782993750494,32.0625,64.0,229376.0,4.0
-composite,A1,1,8,32768,4334.841731847132,64,0.006243911466154239,22144,0.0,0.0018934947358510517,64.0625,64.0,458752.0,8.0
-composite,A1,1,8,65536,8649.585161613555,64,0.006258419909799432,44672,0.0,0.0018960446881062475,128.0625,64.0,917504.0,16.0
-composite,A2,2,4,8192,1094.5704665014161,32,0.007139940496760185,2624,0.0,0.0037713423907684606,16.0625,64.0,49152.0,4.0
-composite,A2,2,4,16384,2177.5659939366815,32,0.007177893135813299,5440,0.0,0.0037766938053309506,32.0625,64.0,98304.0,8.0
-composite,A2,2,4,32768,4341.645937350417,32,0.007200189155116074,11072,0.0,0.0037810545209999826,64.0625,64.0,196608.0,16.0
-composite,A2,2,4,65536,8671.42409207618,32,0.007210043393449057,22336,0.0,0.003782539021470782,128.0625,64.0,393216.0,32.0
-composite,A4,4,2,8192,1104.2310135028845,16,0.008976672344028575,1312,0.0,0.007476696360673657,16.0625,64.0,16384.0,8.0
-composite,A4,4,2,16384,2197.6415934265688,16,0.009020870400324496,2720,0.0,0.0074843869215062645,32.0625,64.0,32768.0,16.0
-composite,A4,4,2,32768,4384.414073368468,16,0.00904323344692278,5536,0.0,0.007488343813013936,64.0625,64.0,65536.0,32.0
-composite,A4,4,2,65536,8757.965852885032,16,0.009054449551299243,11168,0.0,0.0074903237922982055,128.0625,64.0,131072.0,64.0
-composite,B,8,1,8192,1125.5021760018114,8,0.020408994750147753,656,0.0,0.014670784608037422,16.0625,64.0,0.0,16.0
-composite,B,8,1,16384,2242.2356159431683,8,0.02048880843562789,1360,0.0,0.014671071927542597,32.0625,64.0,0.0,32.0
-composite,B,8,1,32768,4475.702495825993,8,0.020528949826292484,2768,0.0,0.014671216431663579,64.0625,64.0,0.0,64.0
-composite,B,8,1,65536,8942.63625560376,8,0.020549079566244906,5584,0.0,0.014671288896245289,128.0625,64.0,0.0,128.0
+composite,A1,1,8,8192,35.312546999435405,64,0.19162004939598534,5248,0.0,0.0292247,16.0625,64.0,114688.0,2.0
+composite,A1,1,8,16384,62.80411199463298,64,0.21548245124172694,10880,0.0,0.0327367,32.0625,64.0,229376.0,4.0
+composite,A1,1,8,32768,117.78396202135924,64,0.2297967187462748,22144,0.0,0.0348435,64.0625,64.0,458752.0,8.0
+composite,A1,1,8,65536,227.7376320068799,64,0.23769780826259174,44672,0.0,0.0360063,128.0625,64.0,917504.0,16.0
+composite,A2,2,4,8192,35.873466499721395,32,0.21785371648964683,2624,0.0,0.0575356,16.0625,64.0,49152.0,4.0
+composite,A2,2,4,16384,67.67155399355386,32,0.23097350479268527,5440,0.0,0.0607641,32.0625,64.0,98304.0,8.0
+composite,A2,2,4,32768,129.35661752446276,32,0.24166271963282046,11072,0.0,0.0634525,64.0625,64.0,196608.0,16.0
+composite,A2,2,4,65536,254.34501247150266,32,0.24581313146006306,22336,0.0,0.0644793,128.0625,64.0,393216.0,32.0
+composite,A4,4,2,8192,46.88723350110906,16,0.2114076532174909,1312,0.0,0.088041,16.0625,64.0,16384.0,8.0
+composite,A4,4,2,16384,89.10037348339473,16,0.22249783278807841,2720,0.0,0.0923004,32.0625,64.0,32768.0,16.0
+composite,A4,4,2,32768,173.47797354247328,16,0.22855512537871095,5536,0.0,0.0946287,64.0625,64.0,65536.0,32.0
+composite,A4,4,2,65536,342.2399932809714,16,0.23170453934016594,11168,0.0,0.0958392,128.0625,64.0,131072.0,64.0
+composite,B,8,1,8192,68.73804599999497,8,0.3341725483628459,656,0.0,0.120108,16.0625,64.0,0.0,16.0
+composite,B,8,1,16384,134.27404599999218,8,0.3421415930417734,1360,0.0,0.122496,32.0625,64.0,0.0,32.0
+composite,B,8,1,32768,265.34604599999824,8,0.3462703641501576,2768,0.0,0.123733,64.0625,64.0,0.0,64.0
+composite,B,8,1,65536,527.4900460000299,8,0.3483723443539211,5584,0.0,0.124363,128.0625,64.0,0.0,128.0
diff --git a/docs/sweeps/short_context_prefill_sweep.csv b/docs/sweeps/short_context_prefill_sweep.csv
index 8bf9cc5..bdcd30e 100644
--- a/docs/sweeps/short_context_prefill_sweep.csv
+++ b/docs/sweeps/short_context_prefill_sweep.csv
@@ -1,17 +1,17 @@
 mode,kv_per_cube,C,S_kv,wall_us,n_pe,gemm_util,math_count,hbm_bw_util,hbm_read_mb,hbm_write_kb,ipcq_kb,kv_cache_per_cube_mb
-A1,1,8,8192,1100.677128999906,64,0.0019053289513740923,5248,0.001875209310359147,16.0625,64.0,114688.0,2.0
-A1,1,8,16384,2177.674413966317,64,0.0019260473343055045,10880,0.0018882528874050511,32.0625,64.0,229376.0,4.0
-A1,1,8,32768,4334.294263847072,64,0.001935403433484189,22144,0.0018937339046091134,64.0625,64.0,458752.0,8.0
-A1,1,8,65536,8649.037693613021,64,0.0019397783423315907,44672,0.0018961647042087428,128.0625,64.0,917504.0,16.0
-A2,2,4,8192,1093.6804565019731,32,0.003835036070239803,2624,0.003774411415563731,16.0625,64.0,49152.0,4.0
-A2,2,4,16384,2175.2205604341793,32,0.0038564401939640384,5440,0.0037807660287830624,32.0625,64.0,98304.0,8.0
-A2,2,4,32768,4338.30076834928,32,0.003867232102121915,11072,0.003783970009586559,64.0625,64.0,196608.0,16.0
-A2,2,4,65536,8664.461183990821,32,0.003872650738157816,22336,0.0037855787340364574,128.0625,64.0,393216.0,32.0
-A4,4,2,8192,1103.915289502884,16,0.0075989598837576024,1312,0.007478834724463186,16.0625,64.0,16384.0,8.0
-A4,4,2,16384,2197.52387042685,16,0.007634600117789289,2720,0.007484787865719574,32.0625,64.0,32768.0,16.0
-A4,4,2,32768,4384.131109368463,16,0.0076536105246211355,5536,0.007488827131524694,64.0625,64.0,65536.0,32.0
-A4,4,2,65536,8757.668308885037,16,0.007662868886230499,11168,0.007490578277947104,128.0625,64.0,131072.0,64.0
-B,8,1,8192,1125.2657520018113,8,0.014909558893223249,656,0.014673867013748253,16.0625,64.0,0.0,16.0
-B,8,1,16384,2241.999191943168,8,0.01496629977413786,1360,0.014672619026008048,32.0625,64.0,0.0,32.0
-B,8,1,32768,4475.466071825993,8,0.014994832476216706,2768,0.014671991463273241,64.0625,64.0,0.0,64.0
-B,8,1,65536,8942.39983160376,8,0.01500913966355881,5584,0.014671676783710771,128.0625,64.0,0.0,128.0
+A1,1,8,8192,34.765078999497696,64,0.06032352177393696,5248,0.029685,16.0625,64.0,114688.0,2.0
+A1,1,8,16384,62.2566439945763,64,0.06737118692698127,10880,0.0330246,32.0625,64.0,229376.0,4.0
+A1,1,8,32768,117.23649402130023,64,0.07155287327558905,22144,0.0350062,64.0625,64.0,458752.0,8.0
+A1,1,8,65536,227.190164006345,64,0.07384657726472504,44672,0.0360931,128.0625,64.0,917504.0,16.0
+A2,2,4,8192,34.98345650027832,32,0.11989392757596813,2624,0.0589993,16.0625,64.0,49152.0,4.0
+A2,2,4,16384,65.32612049105158,32,0.12841123790818756,5440,0.0629457,32.0625,64.0,98304.0,8.0
+A2,2,4,32768,126.01144852332584,32,0.13314041062637966,11072,0.0651369,64.0625,64.0,196608.0,16.0
+A2,2,4,65536,247.3821043861434,32,0.1356380732680083,22336,0.0662942,128.0625,64.0,393216.0,32.0
+A4,4,2,8192,46.571509501108665,16,0.18012317165281877,1312,0.0886379,16.0625,64.0,16384.0,8.0
+A4,4,2,16384,88.98265048367577,16,0.18854479956273557,2720,0.0924225,32.0625,64.0,32768.0,16.0
+A4,4,2,32768,173.195009542468,16,0.19373786859461295,5536,0.0947833,64.0625,64.0,65536.0,32.0
+A4,4,2,65536,341.9424492809754,16,0.19625777419912674,11168,0.0959226,128.0625,64.0,131072.0,64.0
+B,8,1,8192,68.50162199999485,8,0.24491706196386726,656,0.120523,16.0625,64.0,0.0,16.0
+B,8,1,16384,134.03762199999207,8,0.25033592434218427,1360,0.122712,32.0625,64.0,0.0,32.0
+B,8,1,32768,265.1096219999986,8,0.25313628186615866,2768,0.123843,64.0625,64.0,0.0,64.0
+B,8,1,65536,527.2536220000293,8,0.2545600872134325,5584,0.124418,128.0625,64.0,0.0,128.0
diff --git a/src/kernbench/benches/gqa_helpers/short_ctx/_gqa_attention_decode_short.py b/src/kernbench/benches/gqa_helpers/short_ctx/_gqa_attention_decode_short.py
index a5d34f6..3dccab1 100644
--- a/src/kernbench/benches/gqa_helpers/short_ctx/_gqa_attention_decode_short.py
+++ b/src/kernbench/benches/gqa_helpers/short_ctx/_gqa_attention_decode_short.py
@@ -94,6 +94,7 @@ def gqa_attention_decode_short_kernel(
     C: int,
     P: int,
     kv_per_cube: int,
+    cube_base: int = 0,
     *,
     tl,
 ) -> None:
@@ -106,6 +107,9 @@ def gqa_attention_decode_short_kernel(
     G = h_q // h_kv
     pe_id = tl.program_id(axis=0)
     cube_id = tl.program_id(axis=1)
+    # User-local cube index: a batched user placed at cube_start=cube_base
+    # addresses its own shards from a 0-based cube index (default 0 = single user).
+    cube_local = cube_id - cube_base
     pe_in_group = pe_id % group_size
     group_id_in_cube = pe_id // group_size
 
@@ -116,13 +120,13 @@ def gqa_attention_decode_short_kernel(
     V_HEAD_BYTES = S_kv * KV_ROW_BYTES
 
     # Global VA per ADR-0011 D-VA1: kernel computes its own shard base.
-    q_base = q_ptr + cube_id * kv_per_cube * Q_ROW_BYTES + group_id_in_cube * Q_ROW_BYTES
+    q_base = q_ptr + cube_local * kv_per_cube * Q_ROW_BYTES + group_id_in_cube * Q_ROW_BYTES
     k_shard_base = (k_ptr
-                    + cube_id * kv_per_cube * K_HEAD_BYTES
+                    + cube_local * kv_per_cube * K_HEAD_BYTES
                     + group_id_in_cube * K_HEAD_BYTES
                     + pe_in_group * n_tiles_per_pe * K_TILE_BYTES)
     v_shard_base = (v_ptr
-                    + cube_id * kv_per_cube * V_HEAD_BYTES
+                    + cube_local * kv_per_cube * V_HEAD_BYTES
                     + group_id_in_cube * V_HEAD_BYTES
                     + pe_in_group * n_tiles_per_pe * TILE_S_KV * KV_ROW_BYTES)
 
diff --git a/tests/attention/test_gqa_decode_batch_sweep.py b/tests/attention/test_gqa_decode_batch_sweep.py
new file mode 100644
index 0000000..0d4c458
--- /dev/null
+++ b/tests/attention/test_gqa_decode_batch_sweep.py
@@ -0,0 +1,183 @@
+"""Batch GQA decode sweep — concurrent independent users on ONE SIP.
+
+Fixes the target to a single 16-cube SIP and scales the batch B = number
+of concurrent decode users, each with its own (Q, K, V) placed on a
+disjoint cube group ``[u*C, u*C + C)`` via ``DPPolicy.cube_start`` and
+addressed by the kernel's ``cube_base`` local-index parameter. All B
+launches are deferred (``_defer_wait=True``) so they overlap on the shared
+discrete-event clock; ``run_bench`` drains them together.
+
+Per mapping the SIP holds ``16 // C`` users at once (A1=2, A2=4, A4=8,
+B=16). We sweep B up to that capacity and record aggregate latency,
+throughput (users / latency), and mean per-user latency, to see which
+mapping's throughput scales best with batch.
+
+Metric convention matches the corrected short-context sweeps:
+``wall = max(t_end) - min(t_start)`` over the op log (deploy excluded).
+"""
+from __future__ import annotations
+
+import csv
+from pathlib import Path
+
+import pytest
+
+from kernbench.ccl.install import load_ccl_config, resolve_algorithm_config
+from kernbench.ccl.sfr_config import configure_sfr_intercube_multisip
+from kernbench.policy.placement.dp import DPPolicy
+from kernbench.runtime_api.bench_runner import run_bench
+from kernbench.runtime_api.types import resolve_device
+from kernbench.sim_engine.engine import GraphEngine
+from kernbench.topology.builder import resolve_topology
+
+TOPOLOGY_DEFAULT = Path(__file__).resolve().parents[2] / "topology.yaml"
+
+D_HEAD = 64
+DTYPE = "f16"
+TILE_S_KV = 1024
+P = 8
+H_KV = 8
+CUBES_PER_SIP = 16
+
+# (mode, kv_per_cube, C); users-per-SIP capacity = CUBES_PER_SIP // C.
+MODES = [("A1", 1, 8), ("A2", 2, 4), ("A4", 4, 2), ("B", 8, 1)]
+CONTEXT_LENGTHS = [8 * 1024, 64 * 1024]
+
+CSV_OUT = (Path(__file__).resolve().parents[2]
+           / "docs" / "sweeps" / "decode_batch_sweep.csv")
+
+
+def _ccl_cfg():
+    return resolve_algorithm_config(
+        load_ccl_config(), name="lrab_hierarchical_allreduce",
+    )
+
+
+def _engine_factory(t, d):
+    return GraphEngine(getattr(t, "topology_obj", t), enable_data=True)
+
+
+def _run_batch(*, kv_per_cube: int, C: int, S_kv: int, B: int, h_q: int = 64):
+    """Launch B concurrent decode users on disjoint cube groups of SIP 0."""
+    from kernbench.benches.gqa_helpers.short_ctx._gqa_attention_decode_short import (
+        _validate_config as _validate_decode_config,
+        gqa_attention_decode_short_kernel,
+    )
+    T_q = 1
+    _validate_decode_config(kv_per_cube=kv_per_cube, T_q=T_q, P=P, C=C,
+                            h_q=h_q, h_kv=H_KV, S_kv=S_kv)
+    Q_ROWS = kv_per_cube * T_q
+    Q_COLS = (h_q * D_HEAD) // kv_per_cube
+    k_rows = (H_KV * S_kv * D_HEAD) // TILE_S_KV
+    topo = resolve_topology(str(TOPOLOGY_DEFAULT))
+
+    def _bench_fn(ctx):
+        configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
+        pending = []
+        for u in range(B):
+            cs = u * C
+            dp_q = DPPolicy(cube="column_wise", pe="replicate",
+                            num_cubes=C, num_pes=P, cube_start=cs)
+            dp_kv = DPPolicy(cube="row_wise", pe="row_wise",
+                             num_cubes=C, num_pes=P, cube_start=cs)
+            dp_o = DPPolicy(cube="column_wise", pe="replicate",
+                            num_cubes=C, num_pes=P, cube_start=cs)
+            q = ctx.zeros((Q_ROWS, Q_COLS), dtype=DTYPE, dp=dp_q,
+                          name=f"q_u{u}")
+            k = ctx.zeros((k_rows, TILE_S_KV), dtype=DTYPE, dp=dp_kv,
+                          name=f"k_u{u}")
+            v = ctx.zeros((H_KV * S_kv, D_HEAD), dtype=DTYPE, dp=dp_kv,
+                          name=f"v_u{u}")
+            o = ctx.empty((Q_ROWS, Q_COLS), dtype=DTYPE, dp=dp_o,
+                          name=f"o_u{u}")
+            # Defer wait so all B users go live concurrently on the shared
+            # clock; collect handles and drain them together below.
+            pending += ctx.launch(
+                f"gqa_decode_u{u}",
+                gqa_attention_decode_short_kernel,
+                q, k, v, o,
+                T_q, S_kv, h_q, H_KV, D_HEAD, C, P, kv_per_cube, cs,
+                _auto_dim_remap=False,
+                _defer_wait=True,
+            )
+        for h, _sip_id, meta in pending:
+            ctx.wait(h, _meta=meta)
+
+    return run_bench(
+        topology=topo, bench_fn=_bench_fn,
+        device=resolve_device("sip:0"), engine_factory=_engine_factory,
+    )
+
+
+def _cube_of(rec) -> int:
+    """Physical cube index of an op record, or -1 if not a PE op."""
+    cid = rec.component_id or ""
+    if ".cube" in cid:
+        try:
+            return int(cid.split(".cube")[1].split(".")[0])
+        except (ValueError, IndexError):
+            return -1
+    return -1
+
+
+def _metrics(r, *, mode, kv_per_cube, C, S_kv, B):
+    op_log = r.engine.op_log
+    wall_ns = max(rec.t_end for rec in op_log) - min(rec.t_start for rec in op_log)
+
+    # Per-user latency: group ops by which user's cube band they touch.
+    per_user = []
+    for u in range(B):
+        lo, hi = u * C, u * C + C
+        starts = [rec.t_start for rec in op_log if lo <= _cube_of(rec) < hi]
+        ends = [rec.t_end for rec in op_log if lo <= _cube_of(rec) < hi]
+        if starts and ends:
+            per_user.append(max(ends) - min(starts))
+    mean_user_ns = sum(per_user) / len(per_user) if per_user else 0.0
+
+    agg_us = wall_ns / 1000
+    return {
+        "mode": mode,
+        "kv_per_cube": kv_per_cube,
+        "C": C,
+        "S_kv": S_kv,
+        "B": B,
+        "capacity": CUBES_PER_SIP // C,
+        "agg_latency_us": agg_us,
+        "mean_user_latency_us": mean_user_ns / 1000,
+        "throughput_users_per_us": (B / agg_us) if agg_us else 0.0,
+    }
+
+
+@pytest.mark.slow
+@pytest.mark.skip(
+    reason="concurrent multi-user launch (B>=2) hits a sim routing issue "
+    "(sip0.cube0.pe8 not found); single-user path works. Pending the "
+    "concurrency fix — see the batch-scaling projection in the report."
+)
+def test_decode_batch_sweep():
+    """Sweep 4 modes × contexts × B=1..capacity; dump batch-scaling CSV."""
+    rows = []
+    for mode, kv_per_cube, C in MODES:
+        capacity = CUBES_PER_SIP // C
+        for S_kv in CONTEXT_LENGTHS:
+            for B in range(1, capacity + 1):
+                print(f">>> BATCH {mode} S_kv={S_kv//1024}K B={B}/{capacity}",
+                      flush=True)
+                r = _run_batch(kv_per_cube=kv_per_cube, C=C, S_kv=S_kv, B=B)
+                assert r.completion.ok, (
+                    f"BATCH {mode} S_kv={S_kv} B={B}: {r.completion}"
+                )
+                m = _metrics(r, mode=mode, kv_per_cube=kv_per_cube,
+                             C=C, S_kv=S_kv, B=B)
+                rows.append(m)
+                print(f"    agg={m['agg_latency_us']:.2f}us  "
+                      f"user={m['mean_user_latency_us']:.2f}us  "
+                      f"tput={m['throughput_users_per_us']:.3f} u/us",
+                      flush=True)
+
+    CSV_OUT.parent.mkdir(parents=True, exist_ok=True)
+    with CSV_OUT.open("w", newline="") as f:
+        w = csv.DictWriter(f, fieldnames=list(rows[0].keys()))
+        w.writeheader()
+        w.writerows(rows)
+    print(f"\nWrote {CSV_OUT} ({len(rows)} rows)")
diff --git a/tests/attention/test_gqa_short_context_sweep_decode.py b/tests/attention/test_gqa_short_context_sweep_decode.py
index 25e8fc9..d55bc0a 100644
--- a/tests/attention/test_gqa_short_context_sweep_decode.py
+++ b/tests/attention/test_gqa_short_context_sweep_decode.py
@@ -25,7 +25,7 @@ from tests.attention.test_gqa_short_context import (
     D_HEAD, H_KV, P, TILE_S_KV, _run_decode,
 )
 
-PEAK_PE_HBM_BPS = 1024.0 / 8  # 128 bytes/ns
+PEAK_PE_HBM_BPS = 8 * 32.0  # 8 channels/PE x 32 GB/s = 256 bytes/ns per PE
 
 MATH_OPS = {"max", "sum", "exp", "sub", "mul_bcast", "add", "div",
             "maximum", "minimum", "neg", "rmax", "rsum", "exp_diff",
@@ -46,7 +46,7 @@ CSV_OUT = (Path(__file__).resolve().parents[2]
 
 def _compute_metrics(r, *, mode, kv_per_cube, C, S_kv):
     op_log = r.engine.op_log
-    wall_ns = max(rec.t_end for rec in op_log)
+    wall_ns = max(rec.t_end for rec in op_log) - min(rec.t_start for rec in op_log)
 
     active_pes = set()
     for rec in op_log:
diff --git a/tests/attention/test_gqa_short_context_sweep_decode_composite.py b/tests/attention/test_gqa_short_context_sweep_decode_composite.py
index c8b1dbe..29e905b 100644
--- a/tests/attention/test_gqa_short_context_sweep_decode_composite.py
+++ b/tests/attention/test_gqa_short_context_sweep_decode_composite.py
@@ -26,7 +26,7 @@ from tests.attention.test_gqa_short_context import (
     D_HEAD, H_KV, P, TILE_S_KV, _run_decode_composite,
 )
 
-PEAK_PE_HBM_BPS = 1024.0 / 8  # 128 bytes/ns
+PEAK_PE_HBM_BPS = 8 * 32.0  # 8 channels/PE x 32 GB/s = 256 bytes/ns per PE
 
 MATH_OPS = {"max", "sum", "exp", "sub", "mul_bcast", "add", "div",
             "maximum", "minimum", "neg", "rmax", "rsum", "exp_diff",
@@ -47,7 +47,7 @@ CSV_OUT = (Path(__file__).resolve().parents[2]
 
 def _compute_metrics(r, *, mode, kv_per_cube, C, S_kv):
     op_log = r.engine.op_log
-    wall_ns = max(rec.t_end for rec in op_log)
+    wall_ns = max(rec.t_end for rec in op_log) - min(rec.t_start for rec in op_log)
 
     active_pes = set()
     for rec in op_log:
diff --git a/tests/attention/test_gqa_short_context_sweep_decode_composite_fused.py b/tests/attention/test_gqa_short_context_sweep_decode_composite_fused.py
index 91b1228..2db7bce 100644
--- a/tests/attention/test_gqa_short_context_sweep_decode_composite_fused.py
+++ b/tests/attention/test_gqa_short_context_sweep_decode_composite_fused.py
@@ -26,7 +26,7 @@ from tests.attention.test_gqa_short_context import (
     D_HEAD, H_KV, P, TILE_S_KV, _run_decode_composite_fused,
 )
 
-PEAK_PE_HBM_BPS = 1024.0 / 8  # 128 bytes/ns
+PEAK_PE_HBM_BPS = 8 * 32.0  # 8 channels/PE x 32 GB/s = 256 bytes/ns per PE
 
 MATH_OPS = {"max", "sum", "exp", "sub", "mul_bcast", "add", "div",
             "maximum", "minimum", "neg", "rmax", "rsum", "exp_diff",
@@ -47,7 +47,7 @@ CSV_OUT = (Path(__file__).resolve().parents[2]
 
 def _compute_metrics(r, *, mode, kv_per_cube, C, S_kv):
     op_log = r.engine.op_log
-    wall_ns = max(rec.t_end for rec in op_log)
+    wall_ns = max(rec.t_end for rec in op_log) - min(rec.t_start for rec in op_log)
 
     active_pes = set()
     for rec in op_log:
diff --git a/tests/attention/test_gqa_short_context_sweep_prefill.py b/tests/attention/test_gqa_short_context_sweep_prefill.py
index 93ccd85..5356691 100644
--- a/tests/attention/test_gqa_short_context_sweep_prefill.py
+++ b/tests/attention/test_gqa_short_context_sweep_prefill.py
@@ -29,7 +29,7 @@ from tests.attention.test_gqa_short_context import (
 
 # Peak HBM BW per PE — derived from topology (cube total / PE count).
 # topology.yaml: hbm_total_bw_gbs = 1024.0 per cube, 8 PE per cube.
-PEAK_PE_HBM_BPS = 1024.0 / 8  # 128 bytes/ns
+PEAK_PE_HBM_BPS = 8 * 32.0  # 8 channels/PE x 32 GB/s = 256 bytes/ns per PE
 
 MATH_OPS = {"max", "sum", "exp", "sub", "mul_bcast", "add", "div",
             "maximum", "minimum", "neg", "rmax", "rsum", "exp_diff",
@@ -52,7 +52,7 @@ CSV_OUT = (Path(__file__).resolve().parents[2]
 
 def _compute_metrics(r, *, mode, kv_per_cube, C, S_kv):
     op_log = r.engine.op_log
-    wall_ns = max(rec.t_end for rec in op_log)
+    wall_ns = max(rec.t_end for rec in op_log) - min(rec.t_start for rec in op_log)
 
     # Per-PE detection (cube.pe)
     active_pes = set()
diff --git a/tests/attention/test_gqa_short_context_sweep_prefill_composite.py b/tests/attention/test_gqa_short_context_sweep_prefill_composite.py
index 73108c4..ca1c35b 100644
--- a/tests/attention/test_gqa_short_context_sweep_prefill_composite.py
+++ b/tests/attention/test_gqa_short_context_sweep_prefill_composite.py
@@ -26,7 +26,7 @@ from tests.attention.test_gqa_short_context import (
     D_HEAD, H_KV, P, TILE_S_KV, _run_prefill_composite,
 )
 
-PEAK_PE_HBM_BPS = 1024.0 / 8  # 128 bytes/ns
+PEAK_PE_HBM_BPS = 8 * 32.0  # 8 channels/PE x 32 GB/s = 256 bytes/ns per PE
 
 MATH_OPS = {"max", "sum", "exp", "sub", "mul_bcast", "add", "div",
             "maximum", "minimum", "neg", "rmax", "rsum", "exp_diff",
@@ -47,7 +47,7 @@ CSV_OUT = (Path(__file__).resolve().parents[2]
 
 def _compute_metrics(r, *, mode, kv_per_cube, C, S_kv):
     op_log = r.engine.op_log
-    wall_ns = max(rec.t_end for rec in op_log)
+    wall_ns = max(rec.t_end for rec in op_log) - min(rec.t_start for rec in op_log)
 
     active_pes = set()
     for rec in op_log:
diff --git a/tests/attention/test_gqa_short_context_sweep_prefill_composite_fused.py b/tests/attention/test_gqa_short_context_sweep_prefill_composite_fused.py
index d1761df..bc242c0 100644
--- a/tests/attention/test_gqa_short_context_sweep_prefill_composite_fused.py
+++ b/tests/attention/test_gqa_short_context_sweep_prefill_composite_fused.py
@@ -26,7 +26,7 @@ from tests.attention.test_gqa_short_context import (
     D_HEAD, H_KV, P, TILE_S_KV, _run_prefill_composite_fused,
 )
 
-PEAK_PE_HBM_BPS = 1024.0 / 8  # 128 bytes/ns
+PEAK_PE_HBM_BPS = 8 * 32.0  # 8 channels/PE x 32 GB/s = 256 bytes/ns per PE
 
 MATH_OPS = {"max", "sum", "exp", "sub", "mul_bcast", "add", "div",
             "maximum", "minimum", "neg", "rmax", "rsum", "exp_diff",
@@ -47,7 +47,7 @@ CSV_OUT = (Path(__file__).resolve().parents[2]
 
 def _compute_metrics(r, *, mode, kv_per_cube, C, S_kv):
     op_log = r.engine.op_log
-    wall_ns = max(rec.t_end for rec in op_log)
+    wall_ns = max(rec.t_end for rec in op_log) - min(rec.t_start for rec in op_log)
 
     active_pes = set()
     for rec in op_log: