diff --git a/docs/report/1H-codesign-paper/build/main.pdf b/docs/report/1H-codesign-paper/build/main.pdf index e9eb5b1..6a85222 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_decode_long_ctx_6cases_latency.png b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_latency.png index 7dc9882..00fff81 100644 Binary files a/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_latency.png and b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_latency.png differ diff --git a/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_memory.png b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_memory.png index bc30493..99cf247 100644 Binary files a/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_memory.png and b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_memory.png differ diff --git a/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_parallelism.png b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_parallelism.png index 7c4c228..a753c6a 100644 Binary files a/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_parallelism.png and b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_parallelism.png differ diff --git a/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_traffic.png b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_traffic.png index b95121f..d9dba59 100644 Binary files a/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_traffic.png and b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_traffic.png differ diff --git a/docs/report/1H-codesign-paper/sections/03-gemm.tex b/docs/report/1H-codesign-paper/sections/03-gemm.tex index 5ad3515..dcad1e0 100644 --- a/docs/report/1H-codesign-paper/sections/03-gemm.tex +++ b/docs/report/1H-codesign-paper/sections/03-gemm.tex @@ -172,16 +172,16 @@ kernel of \S\ref{sec:gqa} reaches for next — keeping the right working set on-chip so the composite pipeline lands in the compute-rich regime rather than the BW-bound one. -\subsection{Why composite, and not user-orchestrated async loading?} +\subsection{Why composite, and not kernel-orchestrated async loading?} \label{sec:gemm-vs-async} A reader familiar with double-buffered GEMM kernels on conventional hardware may ask: why a hardware-side composite command at all? Why -isn't the obvious user-level pattern --- async-load each operand, +isn't the obvious kernel-level pattern --- async-load each operand, overlap with compute, accumulate --- sufficient? To answer this concretely we contrast composite against two -user-orchestrated baselines that have access to the same single-op +kernel-orchestrated baselines that have access to the same single-op primitives the platform exposes (\textsf{tl.load} for async DMA into TCM, \textsf{tl.dot} for a single-op GEMM command on TCM-resident operands, \textsf{tl.store} for a DMA write-back). Both baselines @@ -199,7 +199,7 @@ way at the runtime API surface to express ``start computing on tile 0 of $B$ while tile 1 is still in flight.'' Load-of-$B$ and GEMM therefore serialize. -\paragraph{Async-tiled (chunked prefetch).} The user-level workaround is to +\paragraph{Async-tiled (chunked prefetch).} The kernel-level workaround is to split $B$ along $K$ into \textsf{TILE\_K}-sized chunks, issue async \textsf{tl.load}s for those chunks, issue one \textsf{tl.dot} per chunk (each blocking only on its own $b_i$), and accumulate via @@ -227,7 +227,7 @@ $K_{\text{KV}}=4096, d_{\text{head}}=128$ already needs --- past the cap. async-full and queue-all async-tiled are therefore not just slower than composite but \emph{architecturally infeasible} at LLM context length. The depth-2 async-tiled kernel is the only -user-level +kernel-level variant whose peak TCM footprint stays $O(2 \cdot \textsf{TILE\_K} \cdot N)$ regardless of $K$, the same order as composite's per-tile streaming buffer. It is the apples-to-apples @@ -249,7 +249,7 @@ onto $K / \textsf{TILE\_K} = 48$ hardware tiles, so per-tile costs amplify into the largest measurable gap. The work content is identical for all four kernels: $\sim$6.3 M f16 MACs and $\sim$386 KiB of $B$ traffic from HBM, which together require -$\sim$786 ns of GEMM-engine compute and $\sim$750 ns of DMA on a +$\sim$786 ns of GEMM-engine compute and $\sim$781 ns of DMA on a saturated per-PE link. What differs is the number of host commands the same work is decomposed into --- 2 for composite (one \textsf{tl.load(A)} plus one composite), 4 for async-full, and 192 @@ -257,7 +257,7 @@ for either async-tiled variant (48 $A$-loads + 48 $B$-loads + 48 \textsf{tl.dot}s + 47 elementwise adds + 1 store). The engine-pipeline-window throughput tracks that decomposition closely: composite reaches \SI{7.18}{\tera\flop\per\second} (post-overlap -limit, only \SI{12}{\percent} below the \SI{8}{\tera\flop\per\second} +limit, only \SI{10}{\percent} below the \SI{8}{\tera\flop\per\second} per-PE GEMM peak), async-full \SI{3.91}{\tera\flop\per\second} (DMA and compute serialize on a single big dot), and both async-tiled variants $\sim$\SI{2.53}{\tera\flop\per\second} (192 commands' worth of @@ -272,7 +272,7 @@ specific simulator mechanisms. four issuance patterns: composite (one command, scheduler streams per-tile internally), async-full (one \textsf{tl.dot} on fully-loaded $B$), async-tiled with depth-2 double-buffer -(TCM-bounded; the only user-level variant that scales to LLM +(TCM-bounded; the only kernel-level variant that scales to LLM context length), and async-tiled with depth-$\infty$ (all B-tiles queued up front; included as a sanity check that the prefetch depth is \emph{not} what separates the async-tiled kernel from @@ -293,7 +293,7 @@ and all three async kernels beat it.} \end{figure*} \paragraph{Decomposing the gap.} Three structural mechanisms separate -composite from the user-level baselines, and they layer. +composite from the kernel-level baselines, and they layer. \emph{1. Inter-engine token routing happens below the host-side dispatch path.} The composite encodes the full @@ -309,7 +309,7 @@ hand-offs total --- behind a single command from the host's point of view. \emph{2. \textsf{tl.dot} cannot replicate that per-tile pipeline at the -user level.} A single-op GEMM command is handled on the GEMM engine as +kernel level.} A single-op GEMM command is handled on the GEMM engine as a single monolithic compute timeout for the supplied $M{\times}K{\times}N$; there is no internal token loop that would let a streaming DMA of $B[i{+}1]$ overlap with the GEMM of $B[i]$ inside one @@ -388,7 +388,7 @@ the four kernels to combine (a) macro-command dispatch at the host boundary (amortizing the structural CPU cost across all the work a single GEMM does), (b) scheduler-internal per-HW-tile streaming of DMA$\rightleftarrows$compute, and (c) TCM-bounded streaming buffer. -User-orchestrated async kernels can have any two of those, not all +Kernel-orchestrated async kernels can have any two of those, not all three: async-full pays one host dispatch (a) but forfeits per-tile overlap (b) and pins all of $B$ in TCM (c); depth-$\infty$ async-tiled achieves inter-chunk overlap but at $N_K$ host dispatches and diff --git a/docs/report/1H-codesign-paper/sections/04-allreduce.tex b/docs/report/1H-codesign-paper/sections/04-allreduce.tex index bf7c317..b731bc7 100644 --- a/docs/report/1H-codesign-paper/sections/04-allreduce.tex +++ b/docs/report/1H-codesign-paper/sections/04-allreduce.tex @@ -127,6 +127,63 @@ across whatever inter-device topology the configuration specifies, with the IPCQ ring buffer placed in on-PE TCM, PE-local HBM, or cube-shared SRAM---the third knob the results section sweeps. +\subsection{Design alternatives} +\label{sec:ipcq-alternatives} + +PE\_IPCQ is one point in a small space of hardware mechanisms for +moving a short message from one PE to a neighbor and signalling its +arrival. Three established alternatives anchor the space, each the +HW realization of a familiar host-networking idea: a \emph{doorbell + +polling} scheme (the classic MMIO doorbell---write the payload by DMA, +write a doorbell, let the peer poll or take an interrupt); a +\emph{hardware message queue} (HMQ, the NVLink-style descriptor engine +that pushes a queue entry to the peer, with large payloads still +riding a second DMA); and a \emph{completion-queue} design (RDMA-CQ, +the InfiniBand/RoCE pattern where a DMA write auto-posts a completion +entry the peer's CQ polls). PE\_IPCQ is the fourth: a hardware ring +with credit return, splitting the control plane into PE\_IPCQ and the +data plane into PE\_DMA, with head updates riding the payload and tail +updates riding a 16\,B side-channel credit (\S\ref{sec:allreduce}). + +\begin{figure}[t] +\centering +\includegraphics[width=\linewidth]{ipcq_alternatives_architecture_stacked.png} +\caption{Per-send data and control flow for the four PE-to-PE +signalling mechanisms (sender\,$\rightarrow$\,NoC\,$\rightarrow$\,receiver). +Doorbell and RDMA-CQ each issue two fabric transactions (payload then +doorbell / completion) and leave the peer polling or taking an +interrupt; HMQ adds a dedicated descriptor engine but still moves large +payloads on a second DMA; PE\_IPCQ folds head-pointer signalling into +the payload flit train and returns the tail credit on a side channel, +so a send is one MMIO write and a receive is a flip-flop read. This is +a \emph{design schematic}, not a measured comparison.} +\label{fig:ipcq-arch} +\end{figure} + +\begin{figure}[t] +\centering +\includegraphics[width=\linewidth]{ipcq_alternatives_decision_matrix.png} +\caption{Why the ring+credit design was chosen, across five criteria: +single-send latency, whether the host CPU sits on the critical path, +whether the receiver must poll or take a wake-up interrupt, whether the +control and data datapaths are duplicated, and whether the mechanism is +right-sized for single-owner PE-to-PE traffic (rather than a +multi-tenant fabric). PE\_IPCQ is the only design that clears every +criterion. The accompanying per-send step-count tally +($\sim$28 control events for IPCQ versus $\sim$38 for HMQ, $\sim$53 for +RDMA-CQ, and $\sim$56 for doorbell+polling) is an \emph{illustrative} +order-of-magnitude comparator over hand-counted pipeline steps---not a +simulator measurement. The measured, simulator-grounded results follow +in the next subsection.} +\label{fig:ipcq-decision} +\end{figure} + +The qualitative comparison motivates the design but is not a +quantitative claim: the cycle-step tallies above are hand-counted +control events, deliberately separated from the measured latencies that +follow. Everything in the results subsection runs on the PE\_IPCQ +substrate and is simulator-grounded. + \subsection{Results} All measurements in this section run on the PE\_IPCQ substrate @@ -138,8 +195,21 @@ the collective sweep builds its own six-device (six-SIP, $2\times3$) configurations---distinct from the two-SIP default of Table~\ref{tab:hw}---and measures all-reduce latency as a function of payload size for three inter-device topologies: a 1D ring, a 2D mesh -(no wrap), and a 2D torus. Table~\ref{tab:allreduce} and -Figure~\ref{fig:allreduce-cmp} report the result. +(no wrap), and a 2D torus (Figure~\ref{fig:allreduce-topo}). +Table~\ref{tab:allreduce} and Figure~\ref{fig:allreduce-cmp} report the +result. + +\begin{figure}[t] +\centering +\includegraphics[width=\linewidth]{allreduce_topology.png} +\caption{The three six-device ($2\times3$) inter-device topologies the +collective sweep runs over, and the hierarchical local-reduce / +global all-reduce-broadcast schedule mapped onto each: a 1D ring, a 2D +mesh (no wrap-around), and a 2D torus (wrap-around links on both axes). +The torus's wrap links shorten the worst-case reduction path, which is +what the latency sweep below rewards.} +\label{fig:allreduce-topo} +\end{figure} \begin{table}[t] \centering diff --git a/docs/report/1H-codesign-paper/sections/05-gqa.tex b/docs/report/1H-codesign-paper/sections/05-gqa.tex index e165309..cb0edb4 100644 --- a/docs/report/1H-codesign-paper/sections/05-gqa.tex +++ b/docs/report/1H-codesign-paper/sections/05-gqa.tex @@ -27,15 +27,83 @@ the established baseline rather than being re-derived. \subsection{Data Placement Policy} \label{sec:gqa-placement} -% TODO: compare data-placement options that apply across both the -% short- and long-context regimes. Candidate axes: -% - KV cache: per-CUBE shard vs. replicate; per-PE shard vs. replicate -% - Q / W_qkv / W_o weights: static partition across CUBEs and PEs -% - Workspace (m, l, O softmax state): scratch arena placement -% The 4-case taxonomy used for long-context decode in -% \S\ref{sec:gqa-long} (Cube-{SP,Repl} x PE-{TP,SP}) is one instantiation -% of this framework; the short-context mapping in \S\ref{sec:gqa-short} -% is another. +Long-context decode is bound by the KV cache, so the first-order design +question is how to place that cache---and the running softmax state it +feeds---across the two hardware axes the machine exposes: the CUBEs and, +within each CUBE, the PEs (here $C{=}8$ CUBEs $\times$ $P{=}8$ PEs, for +$C\!\cdot\!P{=}64$ attention engines over one KV-head group on the +LLaMA-3.1-70B target). Each axis can \emph{replicate} the KV cache or +\emph{shard} it, and a shard can run along the sequence dimension +$S_{kv}$ or the head dimension $d_{\text{head}}$. The cross product is a +small, enumerable taxonomy; six placements span its meaningful corners +(Figure~\ref{fig:gqa-kv-sharding}). + +\begin{figure}[t] +\centering +\includegraphics[width=\linewidth]{gqa_long_ctx_6cases_kv_sharding_diagram.png} +\caption{The six KV-placement strategies, drawn on the +$S_{kv}\!\times\!d_{\text{head}}$ KV tensor (rows = sequence, columns = +head dimension). Cube colour bands and dashed PE dividers show which +axis each level shards. Cases~1--3 either replicate the cache or shard +it on a single axis (8-way at most); Cases~4--6 reach a full 64-way +split, three different ways: Case~4 splits $S_{kv}$ across CUBEs and +$d_{\text{head}}$ across PEs, Case~5 the mirror, and Case~6~$\star$ +splits $S_{kv}$ on \emph{both} axes.} +\label{fig:gqa-kv-sharding} +\end{figure} + +Two quantities decide which placement is viable, and they pull against +each other. The first is \textbf{per-PE KV memory}. With a per-PE HBM +budget of \SI{6.0}{\giga\byte} and \SI{1.76}{\giga\byte} of attention +weights resident, the KV headroom is \SI{4.24}{\giga\byte} per PE. At a +production context of $S_{kv}{=}1\,\text{M}$ tokens the unsharded cache +is \SI{40}{\giga\byte}/PE (Case~1), an 8-way shard is \SI{5}{\giga\byte} +(Cases~2--3)---both \emph{over} the headroom---while only the 64-way +placements bring it to \SI{640}{\mega\byte}/PE (Cases~4--6), comfortably +inside budget. Memory alone therefore eliminates Cases~1--3 at long +context. The second quantity is \textbf{communication per token}, and it +is what separates the three survivors. Sharding $d_{\text{head}}$ +(Cases~4--5) makes each PE hold only a slice of every head, so the +$Q\!\cdot\!K^{\top}$ score is \emph{partial} and must be all-reduced +across the slice owners on every token---a reduction whose volume scales +with $S_{kv}$ ($\sim$\SI{166}{\mega\byte}/token analytically, intra-CUBE +on the NoC for Case~4, inter-CUBE on UCIe for Case~5). Case~6~$\star$ +instead shards $S_{kv}$ on both axes, so every PE computes a +\emph{complete} score over its own token range and only the small running +softmax state $(m,\ell,O)$ is merged across PEs +($\sim$\SI{6.2}{\mega\byte}/token)---a $\sim$27$\times$ lighter collective +than the $d_{\text{head}}$-split designs. + +\begin{figure}[t] +\centering +\includegraphics[width=\linewidth]{gqa_long_ctx_6cases_summary.png} +\caption{Long-context placement analysis at $S_{kv}{=}1\,\text{M}$ tokens. +\emph{Left:} the per-PE HBM budget---\SI{1.76}{\giga\byte} of attention +weights leave \SI{4.24}{\giga\byte} of KV headroom (red line). +\emph{Middle:} per-PE KV memory per case (log scale); only the 64-way +placements (Cases~4--6, \SI{640}{\mega\byte}) clear the headroom, while +the unsharded (\SI{40}{\giga\byte}) and 8-way (\SI{5}{\giga\byte}) cases +overflow. \emph{Right:} analytical communication per token (log scale); +the $d_{\text{head}}$-split Cases~4--5 pay a partial-score all-reduce +($\sim$\SI{166}{\mega\byte}/token) that the both-axes-$S_{kv}$ split of +Case~6~$\star$ avoids ($\sim$\SI{6.2}{\mega\byte}/token, merging only the +softmax state). On these two axes Case~6 (marked $\star$ in the figure) +is the single placement that lands both inside the memory budget and at +low per-token communication; whether that combination is the right one to +pick is a regime-specific question taken up next.} +\label{fig:gqa-budget} +\end{figure} + +These two costs---per-PE KV memory and per-token communication---are +intrinsic properties of each placement, fixed by how it shards the cache +and independent of the workload regime. They do not by themselves name a +winner: a short prompt where the whole cache fits on one PE values low +communication and tolerates replication, whereas a million-token decode +is bound by the memory wall and will pay communication to escape it. +Which placement is appropriate is therefore a per-regime question, which +the short- and long-context subsections that follow answer by running the +options on the simulator and reading off latency, traffic, and the +redundant compute each one induces. \subsection{Inference with Short-Context Length} \label{sec:gqa-short} @@ -75,78 +143,83 @@ measured path; results below reflect the implemented kernel only. % TODO: prefill long-context performance figure. -The four headline panels above stress the kernel at moderate context -lengths. Long-context decode---the regime where KV cache size, not -attention compute, sets serving cost---turns the choice of how to -parallelize across cubes and PEs into a first-order design knob. We -compare four strategies on the LLaMA-3.1-70B single-KV-head-group -target (8 CUBEs $\times$ 8 PEs, one KV-head group): - -\begin{itemize}\setlength\itemsep{1pt} - \item \textbf{Case 1} (Cube-SP $\times$ PE-TP): KV split by $S_{kv}$ - across CUBEs; PEs tensor-parallel on the batch dimension - (wastes PE-TP work at $B{=}1$). - \item \textbf{Case 2} (Cube-Repl $\times$ PE-TP): full KV - replicated to every CUBE; PEs tensor-parallel on batch. - \item \textbf{Case 3} (Cube-Repl $\times$ PE-SP): full KV - replicated; PEs sequence-parallel on $S_{kv}$ with an - intra-CUBE all-reduce. - \item \textbf{Case 4} ($\star$, Cube-SP $\times$ PE-SP): KV split - 64-way (across both CUBEs and PEs) with a two-phase all-reduce - on the running softmax state $(m, \ell, O)$. -\end{itemize} +Long-context decode is the regime where the KV cache, not attention +compute, sets serving cost, so the placement question of +\S\ref{sec:gqa-placement} becomes decisive here. To pick the right +placement for this regime we run each of the six options as a fused +decode kernel on the simulator (one decode step on the LLaMA-3.1-70B +single-KV-head-group target, $C{=}8$ CUBEs $\times$ $P{=}8$ PEs) at a +tractable $S_{kv}{=}8192$, and read off end-to-end latency, on-device op +traffic, and the redundant compute each one induces. The swept context is +small enough that all six fit in memory at $S_{kv}{=}8192$; the +placements the long-context memory budget rules out (Cases~1--3) are +drawn in red, run here only to expose their issue and communication +structure. \begin{figure}[t] \centering -\includegraphics[width=\linewidth]{gqa_decode_long_ctx_4cases_latency.png} -\caption{End-to-end decode latency per parallelism strategy -(LLaMA-3.1-70B single-KV-head group, 8 CUBEs $\times$ 8 PEs). -Replication into CUBEs (Cases 2/3) wins the latency race -(\SI{20.2}{\micro\second} for Case 3), but Case~4 ($\star$, KV split -64-way) finishes within \SI{14}{\micro\second} of the leader while -paying a different cost---visible in Figure~\ref{fig:gqa-4cases-mem}.} -\label{fig:gqa-4cases-lat} +\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_latency.png} +\caption{Measured end-to-end decode latency per placement +($S_{kv}{=}8192$; red = ruled out by the long-context memory budget, +blue = the predicted Pareto choice). The fastest raw latency belongs to +Case~3 (\SI{17.8}{\micro\second})---but Case~3 replicates the full KV +cache into every CUBE, so it overflows the per-PE budget at production +context and wastes 8$\times$ the compute +(Figure~\ref{fig:gqa-6cases-par}). Among the placements that actually fit +1\,M-token memory (Cases~4--6), Case~6~$\star$ is the fastest +(\SI{30.6}{\micro\second}, versus \SI{31.4}{} and \SI{34.5}{\micro\second} +for the $d_{\text{head}}$-split Cases~4 and~5)---making it the placement +of choice for long-context decode.} +\label{fig:gqa-6cases-lat} \end{figure} \begin{figure}[t] \centering -\includegraphics[width=\linewidth]{gqa_decode_long_ctx_4cases_memory.png} -\caption{Per-CUBE KV memory footprint for the four cases. Cases 1 -and 4---both with the KV cache split across cubes (Cube-SP)---hold -only \SI{0.5}{\mebi\byte} of KV state per CUBE; Cases 2 and 3, which -replicate the full KV, hold \SI{4}{\mebi\byte} per CUBE, an -\textbf{8$\times$} blowup at this configuration that scales linearly -with context length.} -\label{fig:gqa-4cases-mem} +\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_parallelism.png} +\caption{Redundant compute per placement, measured as active-PE +$\times$ $S_{\text{local}}$ (PE-tokens; lower means less wasted work). +The minimum is \num{8192} PE-tokens---one pass over the sequence. +Case~3 inflates this 8$\times$ to \num{65536} by replicating the KV +cache across all eight CUBEs so every CUBE redundantly re-attends the +whole sequence; the $d_{\text{head}}$-split Cases~4--5 likewise carry +\num{65536} because each token is processed across eight head slices. +Case~6~$\star$ achieves the full 64-way split at the minimal +\num{8192} PE-tokens---fully parallel, no replication.} +\label{fig:gqa-6cases-par} \end{figure} \begin{figure}[t] \centering -\includegraphics[width=\linewidth]{gqa_decode_long_ctx_4cases_traffic.png} -\caption{Per-case op-count breakdown. The replicated-KV PE-TP design -(Case 2) avoids almost all on-device communication -(\textasciitilde0 IPCQ copies), but at the cost of KV memory. -Case~4's two-phase reduce charges \textasciitilde190 IPCQ copies and -\textasciitilde190 DMA reads---this is the traffic that PE\_IPCQ -(\S\ref{sec:allreduce}) is built to absorb at on-device speed.} -\label{fig:gqa-4cases-traffic} +\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_traffic.png} +\caption{Measured on-device op traffic per placement. The unsharded +Case~1 issues no IPCQ copies (each PE has the full cache, nothing to +reduce); the single-axis Case~3 charges 168. Among the 64-way splits, +the $d_{\text{head}}$-split Cases~4--5 charge the most---280 IPCQ copies +and 8 DMA writes each, the partial-score all-reduce that head-slicing +forces---while Case~6~$\star$ needs only 189 IPCQ copies and a single DMA +write, because it merges just the running softmax state $(m,\ell,O)$ +rather than partial scores. This is the on-device collective traffic that +PE\_IPCQ and the torus links of \S\ref{sec:allreduce} are provisioned to +absorb at link speed.} +\label{fig:gqa-6cases-traffic} \end{figure} -Three things stand out. First, the fastest case in pure latency -(Case~3, \SI{20.2}{\micro\second}) is also the most memory-hungry, -requiring the full KV state on every CUBE---an option that fails to -scale once context length blows past the per-CUBE budget. Second, -Case~4's KV-split design gives back roughly \SI{14}{\micro\second} -versus Case~3 in exchange for an \textbf{8$\times$} KV-memory -reduction; for practical long-context serving where KV capacity is -the binding constraint, this is the trade the design chooses -(marked $\star$). Third, Case~4 pays its way in -\emph{communication}: the op-count panel shows \textasciitilde190 -IPCQ copies and \textasciitilde190 DMA reads, precisely the -on-device collective traffic that PE\_IPCQ and the torus links of -\S\ref{sec:allreduce} are provisioned to move quickly---so the -``slower'' strategy is in fact the one that fully cashes in the -communication-side codesign work of this report. +The measurements select the right placement for this regime. The fastest +raw latency (Case~3, \SI{17.8}{\micro\second}) comes from replicating the +full KV cache into every CUBE---which is exactly the placement the +long-context memory budget forbids, and which the parallelism panel shows +wastes 8$\times$ the compute. Restricting attention to the placements that +fit production-context memory (the 64-way splits, Cases~4--6), the choice +is Case~6~$\star$: it is the fastest of the three +(\SI{30.6}{\micro\second}), and the op-count panel shows why---its +softmax-state-only reduction charges 189 IPCQ copies and one DMA write +against the 280 copies and 8 DMA writes the $d_{\text{head}}$-split +Cases~4--5 pay for their partial-score all-reduce. For long-context +decode, then, the appropriate data placement is the both-axes sequence +shard (Case~6): it is the cheapest-communicating member of the only +memory-feasible family, and the cross-PE softmax reduction it does pay is +precisely the traffic the communication-side codesign of this report is +built to move quickly. \subsection{Comprehensive Analysis} \label{sec:gqa-analysis} diff --git a/src/kernbench/benches/1H_milestone_output/gemm/gemm_hbm_bw_util.png b/src/kernbench/benches/1H_milestone_output/gemm/gemm_hbm_bw_util.png index 5cf6210..1c8ab6b 100644 Binary files a/src/kernbench/benches/1H_milestone_output/gemm/gemm_hbm_bw_util.png and b/src/kernbench/benches/1H_milestone_output/gemm/gemm_hbm_bw_util.png differ diff --git a/src/kernbench/benches/1H_milestone_output/gemm/gemm_mac_utilization_measured.png b/src/kernbench/benches/1H_milestone_output/gemm/gemm_mac_utilization_measured.png index 152272a..05abcfd 100644 Binary files a/src/kernbench/benches/1H_milestone_output/gemm/gemm_mac_utilization_measured.png and b/src/kernbench/benches/1H_milestone_output/gemm/gemm_mac_utilization_measured.png differ diff --git a/src/kernbench/benches/1H_milestone_output/gemm/gemm_mac_utilization_theoretical_vs_measured.png b/src/kernbench/benches/1H_milestone_output/gemm/gemm_mac_utilization_theoretical_vs_measured.png index f9b5683..472c2c6 100644 Binary files a/src/kernbench/benches/1H_milestone_output/gemm/gemm_mac_utilization_theoretical_vs_measured.png and b/src/kernbench/benches/1H_milestone_output/gemm/gemm_mac_utilization_theoretical_vs_measured.png differ diff --git a/src/kernbench/benches/1H_milestone_output/gemm/gemm_per_pe_tflops.png b/src/kernbench/benches/1H_milestone_output/gemm/gemm_per_pe_tflops.png index 866b72b..c3bfe0d 100644 Binary files a/src/kernbench/benches/1H_milestone_output/gemm/gemm_per_pe_tflops.png and b/src/kernbench/benches/1H_milestone_output/gemm/gemm_per_pe_tflops.png differ diff --git a/src/kernbench/benches/1H_milestone_output/gemm/gemm_stage_breakdown.png b/src/kernbench/benches/1H_milestone_output/gemm/gemm_stage_breakdown.png index 781d47e..e28f96e 100644 Binary files a/src/kernbench/benches/1H_milestone_output/gemm/gemm_stage_breakdown.png and b/src/kernbench/benches/1H_milestone_output/gemm/gemm_stage_breakdown.png differ diff --git a/src/kernbench/benches/1H_milestone_output/gemm/gemm_sweep.json b/src/kernbench/benches/1H_milestone_output/gemm/gemm_sweep.json index 6816a0b..cc1b11a 100644 --- a/src/kernbench/benches/1H_milestone_output/gemm/gemm_sweep.json +++ b/src/kernbench/benches/1H_milestone_output/gemm/gemm_sweep.json @@ -33,7 +33,7 @@ "bytes_hbm": 6144, "arith_intensity": 10.666666666666666, "tile_count_expected": 1, - "sim_wall_clock_s": 0.112, + "sim_wall_clock_s": 0.226, "engines": { "pe_dma": { "occupancy_ns": 79.0, @@ -100,7 +100,7 @@ "bytes_hbm": 6144, "arith_intensity": 10.666666666666666, "tile_count_expected": 1, - "sim_wall_clock_s": 0.108, + "sim_wall_clock_s": 0.186, "engines": { "pe_dma": { "occupancy_ns": 71.0, @@ -167,11 +167,11 @@ "bytes_hbm": 6144, "arith_intensity": 10.666666666666666, "tile_count_expected": 1, - "sim_wall_clock_s": 0.105, + "sim_wall_clock_s": 0.181, "engines": { "pe_dma": { "occupancy_ns": 63.0, - "wall_ns": 63.0, + "wall_ns": 51.0, "record_count": 3 }, "pe_fetch_store": { @@ -222,7 +222,7 @@ "record_count": 0 } }, - "pe_window_ns": 183.394, + "pe_window_ns": 151.394, "composite_window_ns": 77.38400000000001 }, { @@ -234,7 +234,7 @@ "bytes_hbm": 10240, "arith_intensity": 12.8, "tile_count_expected": 1, - "sim_wall_clock_s": 0.169, + "sim_wall_clock_s": 0.314, "engines": { "pe_dma": { "occupancy_ns": 79.0, @@ -301,7 +301,7 @@ "bytes_hbm": 10240, "arith_intensity": 12.8, "tile_count_expected": 1, - "sim_wall_clock_s": 0.163, + "sim_wall_clock_s": 0.191, "engines": { "pe_dma": { "occupancy_ns": 79.0, @@ -368,11 +368,11 @@ "bytes_hbm": 10240, "arith_intensity": 12.8, "tile_count_expected": 1, - "sim_wall_clock_s": 0.104, + "sim_wall_clock_s": 0.186, "engines": { "pe_dma": { - "occupancy_ns": 79.0, - "wall_ns": 79.0, + "occupancy_ns": 86.0, + "wall_ns": 66.0, "record_count": 3 }, "pe_fetch_store": { @@ -423,7 +423,7 @@ "record_count": 0 } }, - "pe_window_ns": 191.394, + "pe_window_ns": 166.394, "composite_window_ns": 77.38400000000001 }, { @@ -435,7 +435,7 @@ "bytes_hbm": 18432, "arith_intensity": 14.222222222222221, "tile_count_expected": 2, - "sim_wall_clock_s": 0.162, + "sim_wall_clock_s": 0.318, "engines": { "pe_dma": { "occupancy_ns": 158.995, @@ -502,7 +502,7 @@ "bytes_hbm": 18432, "arith_intensity": 14.222222222222221, "tile_count_expected": 2, - "sim_wall_clock_s": 0.183, + "sim_wall_clock_s": 0.188, "engines": { "pe_dma": { "occupancy_ns": 139.995, @@ -569,11 +569,11 @@ "bytes_hbm": 18432, "arith_intensity": 14.222222222222221, "tile_count_expected": 2, - "sim_wall_clock_s": 0.106, + "sim_wall_clock_s": 0.186, "engines": { "pe_dma": { - "occupancy_ns": 112.0, - "wall_ns": 107.0, + "occupancy_ns": 134.0, + "wall_ns": 98.0, "record_count": 3 }, "pe_fetch_store": { @@ -624,7 +624,7 @@ "record_count": 0 } }, - "pe_window_ns": 223.77800000000002, + "pe_window_ns": 214.77800000000002, "composite_window_ns": 93.76800000000003 }, { @@ -636,7 +636,7 @@ "bytes_hbm": 49152, "arith_intensity": 21.333333333333332, "tile_count_expected": 8, - "sim_wall_clock_s": 0.177, + "sim_wall_clock_s": 0.308, "engines": { "pe_dma": { "occupancy_ns": 1798.8600000000001, @@ -703,7 +703,7 @@ "bytes_hbm": 49152, "arith_intensity": 21.333333333333332, "tile_count_expected": 8, - "sim_wall_clock_s": 0.22, + "sim_wall_clock_s": 0.263, "engines": { "pe_dma": { "occupancy_ns": 819.8600000000001, @@ -770,11 +770,11 @@ "bytes_hbm": 49152, "arith_intensity": 21.333333333333332, "tile_count_expected": 8, - "sim_wall_clock_s": 0.112, + "sim_wall_clock_s": 0.183, "engines": { "pe_dma": { - "occupancy_ns": 283.0, - "wall_ns": 238.07200000000012, + "occupancy_ns": 305.0, + "wall_ns": 229.07200000000012, "record_count": 6 }, "pe_fetch_store": { @@ -825,7 +825,7 @@ "record_count": 0 } }, - "pe_window_ns": 418.0820000000001, + "pe_window_ns": 409.0820000000001, "composite_window_ns": 192.07200000000012 }, { @@ -837,7 +837,7 @@ "bytes_hbm": 395264, "arith_intensity": 15.917098445595855, "tile_count_expected": 48, - "sim_wall_clock_s": 0.2, + "sim_wall_clock_s": 0.293, "engines": { "pe_dma": { "occupancy_ns": 56319.35999999987, @@ -904,7 +904,7 @@ "bytes_hbm": 395264, "arith_intensity": 15.917098445595855, "tile_count_expected": 48, - "sim_wall_clock_s": 0.237, + "sim_wall_clock_s": 0.292, "engines": { "pe_dma": { "occupancy_ns": 20236.359999999866, @@ -912,13 +912,13 @@ "record_count": 50 }, "pe_fetch_store": { - "occupancy_ns": 1543.9999999999998, + "occupancy_ns": 1543.9999999999995, "wall_ns": 791.9999999999998, "record_count": 49 }, "pe_gemm": { - "occupancy_ns": 1219.584000000017, - "wall_ns": 786.4320000000007, + "occupancy_ns": 1219.5840000000169, + "wall_ns": 786.4320000000005, "record_count": 48 }, "pe_math": { @@ -939,7 +939,7 @@ "record_count": 1 }, "FETCH": { - "occupancy_ns": 1535.9999999999998, + "occupancy_ns": 1535.9999999999995, "wall_ns": 783.9999999999998, "record_count": 48 }, @@ -949,8 +949,8 @@ "record_count": 1 }, "GEMM": { - "occupancy_ns": 1219.584000000017, - "wall_ns": 786.4320000000007, + "occupancy_ns": 1219.5840000000169, + "wall_ns": 786.4320000000005, "record_count": 48 }, "MATH": { @@ -971,10 +971,10 @@ "bytes_hbm": 395264, "arith_intensity": 15.917098445595855, "tile_count_expected": 48, - "sim_wall_clock_s": 0.208, + "sim_wall_clock_s": 0.237, "engines": { "pe_dma": { - "occupancy_ns": 2310.0, + "occupancy_ns": 2342.0, "wall_ns": 1569.9999999999998, "record_count": 3 }, @@ -1038,7 +1038,7 @@ "bytes_hbm": 36864, "arith_intensity": 7.111111111111111, "tile_count_expected": 8, - "sim_wall_clock_s": 0.089, + "sim_wall_clock_s": 0.259, "engines": { "pe_dma": { "occupancy_ns": 1798.8600000000001, @@ -1105,7 +1105,7 @@ "bytes_hbm": 36864, "arith_intensity": 7.111111111111111, "tile_count_expected": 8, - "sim_wall_clock_s": 0.155, + "sim_wall_clock_s": 0.313, "engines": { "pe_dma": { "occupancy_ns": 795.8600000000001, @@ -1172,11 +1172,11 @@ "bytes_hbm": 36864, "arith_intensity": 7.111111111111111, "tile_count_expected": 8, - "sim_wall_clock_s": 0.234, + "sim_wall_clock_s": 0.32, "engines": { "pe_dma": { "occupancy_ns": 258.0, - "wall_ns": 218.07200000000012, + "wall_ns": 206.07200000000012, "record_count": 6 }, "pe_fetch_store": { @@ -1227,7 +1227,7 @@ "record_count": 0 } }, - "pe_window_ns": 418.0820000000001, + "pe_window_ns": 386.0820000000001, "composite_window_ns": 192.07200000000012 }, { @@ -1239,7 +1239,7 @@ "bytes_hbm": 36864, "arith_intensity": 7.111111111111111, "tile_count_expected": 16, - "sim_wall_clock_s": 0.119, + "sim_wall_clock_s": 0.304, "engines": { "pe_dma": { "occupancy_ns": 7854.6320000000005, @@ -1306,7 +1306,7 @@ "bytes_hbm": 36864, "arith_intensity": 7.111111111111111, "tile_count_expected": 16, - "sim_wall_clock_s": 0.183, + "sim_wall_clock_s": 0.165, "engines": { "pe_dma": { "occupancy_ns": 3779.6320000000005, @@ -1373,11 +1373,11 @@ "bytes_hbm": 36864, "arith_intensity": 7.111111111111111, "tile_count_expected": 16, - "sim_wall_clock_s": 0.179, + "sim_wall_clock_s": 0.236, "engines": { "pe_dma": { "occupancy_ns": 858.0, - "wall_ns": 183.0, + "wall_ns": 171.0, "record_count": 18 }, "pe_fetch_store": { @@ -1428,7 +1428,7 @@ "record_count": 0 } }, - "pe_window_ns": 511.01, + "pe_window_ns": 479.01, "composite_window_ns": 405.0 }, { @@ -1440,7 +1440,7 @@ "bytes_hbm": 1572864, "arith_intensity": 170.66666666666666, "tile_count_expected": 2048, - "sim_wall_clock_s": 1.241, + "sim_wall_clock_s": 2.164, "engines": { "pe_dma": { "occupancy_ns": 104867178.97599977, @@ -1507,7 +1507,7 @@ "bytes_hbm": 1572864, "arith_intensity": 170.66666666666666, "tile_count_expected": 2048, - "sim_wall_clock_s": 0.742, + "sim_wall_clock_s": 1.398, "engines": { "pe_dma": { "occupancy_ns": 37760375.97599977, @@ -1520,7 +1520,7 @@ "record_count": 2304 }, "pe_gemm": { - "occupancy_ns": 36417.40799998891, + "occupancy_ns": 36417.407999988936, "wall_ns": 33554.43199999761, "record_count": 2048 }, @@ -1547,12 +1547,12 @@ "record_count": 2048 }, "STORE": { - "occupancy_ns": 130151.58400000146, + "occupancy_ns": 130151.58400000144, "wall_ns": 31914.784000000003, "record_count": 256 }, "GEMM": { - "occupancy_ns": 36417.40799998891, + "occupancy_ns": 36417.407999988936, "wall_ns": 33554.43199999761, "record_count": 2048 }, @@ -1574,10 +1574,10 @@ "bytes_hbm": 1572864, "arith_intensity": 170.66666666666666, "tile_count_expected": 2048, - "sim_wall_clock_s": 0.406, + "sim_wall_clock_s": 0.477, "engines": { "pe_dma": { - "occupancy_ns": 142065.0, + "occupancy_ns": 142097.0, "wall_ns": 6170.0, "record_count": 258 }, @@ -1587,8 +1587,8 @@ "record_count": 2304 }, "pe_gemm": { - "occupancy_ns": 838467.5839981049, - "wall_ns": 33554.43199999738, + "occupancy_ns": 838467.5839981119, + "wall_ns": 33554.431999997396, "record_count": 2048 }, "pe_math": { @@ -1619,8 +1619,8 @@ "record_count": 256 }, "GEMM": { - "occupancy_ns": 838467.5839981049, - "wall_ns": 33554.43199999738, + "occupancy_ns": 838467.5839981119, + "wall_ns": 33554.431999997396, "record_count": 2048 }, "MATH": { diff --git a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_latency.png b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_latency.png index 7dc9882..00fff81 100644 Binary files a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_latency.png and b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_latency.png differ diff --git a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_memory.png b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_memory.png index bc30493..99cf247 100644 Binary files a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_memory.png and b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_memory.png differ diff --git a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_parallelism.png b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_parallelism.png index 7c4c228..a753c6a 100644 Binary files a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_parallelism.png and b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_parallelism.png differ diff --git a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_traffic.png b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_traffic.png index b95121f..d9dba59 100644 Binary files a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_traffic.png and b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/gqa_decode_long_ctx_6cases_traffic.png differ diff --git a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/sweep_decode.json b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/sweep_decode.json index 0de23f4..5232a1f 100644 --- a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/sweep_decode.json +++ b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/sweep_decode.json @@ -4,7 +4,9 @@ "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_sp", "single_kv_group_decode_long_ctx_gqa_cube_repl_pe_tp", "single_kv_group_decode_long_ctx_gqa_cube_repl_pe_sp", - "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp" + "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp", + "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp_dhead", + "single_kv_group_decode_long_ctx_gqa_cube_tp_dhead_pe_sp" ], "rows": [ { @@ -23,11 +25,11 @@ "dma_read_count": 192, "dma_write_count": 1 }, - "latency_ns": 34030.00300000079, + "latency_ns": 30646.00300000079, "engine_occupancy_ns": { "pe_gemm": 4194.303999999538, "pe_math": 2282.0, - "pe_dma": 1138201.8199999991, + "pe_dma": 1814643.714999999, "pe_fetch_store": 0, "pe_ipcq": 0, "pe_cpu": 0 @@ -49,11 +51,11 @@ "dma_read_count": 17, "dma_write_count": 1 }, - "latency_ns": 25428.38200085424, + "latency_ns": 21332.36200087145, "engine_occupancy_ns": { "pe_gemm": 4194.303999997675, "pe_math": 714.0, - "pe_dma": 18217.080000881106, + "pe_dma": 26829.080001339316, "pe_fetch_store": 0, "pe_ipcq": 0, "pe_cpu": 0 @@ -75,11 +77,11 @@ "dma_read_count": 192, "dma_write_count": 1 }, - "latency_ns": 20152.09400000231, + "latency_ns": 17826.79400000139, "engine_occupancy_ns": { "pe_gemm": 33554.431999996305, "pe_math": 5684.0, - "pe_dma": 853218.7100000716, + "pe_dma": 1347469.350000088, "pe_fetch_store": 0, "pe_ipcq": 0, "pe_cpu": 0 @@ -101,11 +103,63 @@ "dma_read_count": 24, "dma_write_count": 1 }, - "latency_ns": 26987.291500001098, + "latency_ns": 25750.79150000075, "engine_occupancy_ns": { "pe_gemm": 4194.303999999538, "pe_math": 714.0, - "pe_dma": 126624.99000000354, + "pe_dma": 200429.48000000394, + "pe_fetch_store": 0, + "pe_ipcq": 0, + "pe_cpu": 0 + } + }, + { + "panel": "single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp_dhead", + "kind": "decode_long_ctx_cube_sp_pe_tp_dhead", + "C": 8, + "P": 8, + "T_q": 1, + "S_kv": 8192, + "d_head": 128, + "h_q": 8, + "h_kv": 1, + "op_log_summary": { + "gemm_count": 128, + "ipcq_copy_count": 280, + "dma_read_count": 192, + "dma_write_count": 8 + }, + "latency_ns": 31385.502000000277, + "engine_occupancy_ns": { + "pe_gemm": 4194.303999999545, + "pe_math": 10392.0, + "pe_dma": 1815478.0114999914, + "pe_fetch_store": 0, + "pe_ipcq": 0, + "pe_cpu": 0 + } + }, + { + "panel": "single_kv_group_decode_long_ctx_gqa_cube_tp_dhead_pe_sp", + "kind": "decode_long_ctx_cube_tp_dhead_pe_sp", + "C": 8, + "P": 8, + "T_q": 1, + "S_kv": 8192, + "d_head": 128, + "h_q": 8, + "h_kv": 1, + "op_log_summary": { + "gemm_count": 128, + "ipcq_copy_count": 280, + "dma_read_count": 192, + "dma_write_count": 8 + }, + "latency_ns": 34480.97100000052, + "engine_occupancy_ns": { + "pe_gemm": 4194.303999999538, + "pe_math": 10392.0, + "pe_dma": 2093451.153999999, "pe_fetch_store": 0, "pe_ipcq": 0, "pe_cpu": 0 diff --git a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/sweep_prefill.json b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/sweep_prefill.json index 70989fb..de7047d 100644 --- a/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/sweep_prefill.json +++ b/src/kernbench/benches/1H_milestone_output/gqa/long_ctx/sweep_prefill.json @@ -23,11 +23,11 @@ "dma_read_count": 192, "dma_write_count": 8 }, - "latency_ns": 86702.76799999712, + "latency_ns": 75610.43899999752, "engine_occupancy_ns": { - "pe_gemm": 134217.72799998685, + "pe_gemm": 134217.72799999084, "pe_math": 546656.0, - "pe_dma": 2388650.3709999938, + "pe_dma": 4511934.237999992, "pe_fetch_store": 0, "pe_ipcq": 0, "pe_cpu": 0 @@ -49,11 +49,11 @@ "dma_read_count": 1088, "dma_write_count": 64 }, - "latency_ns": 93564.33200003332, + "latency_ns": 67203.7500000147, "engine_occupancy_ns": { "pe_gemm": 134217.72799998525, - "pe_math": 81279.99999999997, - "pe_dma": 177951.73600022023, + "pe_math": 81280.0, + "pe_dma": 446811.4240000146, "pe_fetch_store": 0, "pe_ipcq": 0, "pe_cpu": 0 @@ -75,11 +75,11 @@ "dma_read_count": 12288, "dma_write_count": 64 }, - "latency_ns": 190579.57100032456, + "latency_ns": 97218.17750020314, "engine_occupancy_ns": { "pe_gemm": 1073741.8240003586, "pe_math": 635904.0, - "pe_dma": 1756854.05751659, + "pe_dma": 3291993.537524393, "pe_fetch_store": 0, "pe_ipcq": 0, "pe_cpu": 0 @@ -101,11 +101,11 @@ "dma_read_count": 192, "dma_write_count": 64 }, - "latency_ns": 38876.32300000038, + "latency_ns": 39045.51900000019, "engine_occupancy_ns": { "pe_gemm": 134217.7280000001, "pe_math": 81280.0, - "pe_dma": 916284.0330000015, + "pe_dma": 1860586.0079999897, "pe_fetch_store": 0, "pe_ipcq": 0, "pe_cpu": 0