T ECHNICAL R EPORT: AI-A SSISTED G ATED D ELTA N ET O PTIMIZATION ON NVIDIA B LACKWELL
arXiv:2607.16831v1 [cs.DC] 18 Jul 2026
Hyunjun Shin 1 Jiseung Jang 1 Jaewoo Maeng 1 Hyunjun Kim 1
A BSTRACT AI-assisted GPU programming is often framed as a kernel-generation loop: ask a model to produce faster CUDA code, benchmark the result, and repeat. This case study argues that contest-grade optimization involves more than improving the kernel body. We examine the Agent-Assisted submission by our team, MSI NFER, to the MLSys 2026 FlashInfer Contest. The submission optimized Gated DeltaNet decode and prefill on NVIDIA B200/Blackwell and achieved an official 1.58× speedup, with approximate average latencies of 9.315 µs for decode and 239.48 µs for prefill. Our experience shows that even effective local kernel improvements can plateau when a workload requires structural reformulation and evaluator-aligned measurement. We therefore characterize AI-assisted kernel optimization as an end-to-end systems problem that encompasses algorithm design, workload specialization, measurement tooling, build and evaluation surfaces, evaluator alignment, and human interpretation. 1
I NTRODUCTION
AI-assisted systems programming promises to shorten the path from an optimization idea to executable GPU code. Large language models can propose CUDA kernels, rewrite indexing logic, generate benchmark scripts, and summarize profiling output. In a contest setting, however, the useful unit of optimization is not the kernel body alone. The submitted artifact must satisfy the evaluator’s build contract, run on the target architecture, pass all correctness workloads, and measure well under the same baseline and timing conditions used for ranking. This paper studies that broader system through the AgentAssisted submission by our team, MSI NFER, to the MLSys 2026 FlashInfer Contest (FlashInfer Contest Organizers, 2026a). The submission targeted Gated DeltaNet (GDN) decode and prefill on NVIDIA B200/Blackwell. Our team produced a correct final submission with an official 1.58× Gated DeltaNet speedup. Its final approximate latencies were 9.315 µs for decode and 239.48 µs for prefill, with all 54 decode and 100 prefill evaluation workloads passing correctness. This report is not a winner’s retrospective. The central observation is instead that our team made meaningful progress through disciplined local search, but the prefill path remained close to a sequential recurrence. That out1 Independent Researcher. Correspondence to: Hyunjun Shin <[email protected]>, Jiseung Jang <[email protected]>, Jaewoo Maeng <[email protected]>, Hyunjun Kim <[email protected]>.
Technical report.
come is useful because it exposes where AI-assisted microoptimization stops being enough: the search loop can improve arithmetic, data movement, launch settings, and measurement discipline, yet still miss a larger change in computational shape. Measurement and algorithm design were coupled throughout the project. A workload-specialized prefill strategy only pays off if the team can reliably tell which workload regimes improved and which regressed. The inspected MSI NFER workflow accumulated Modal benchmark histories, NCU profiles, scorecards, workload inventories, and phase timers, but it did not center a single official-adjacent kernel-latency measurement primitive during most of the search. The result was a broad evidence base that helped reject many local ideas but provided a weaker signal for promoting workloadspecific structural rewrites. The central thesis is: Disciplined measurement and AI-assisted microoptimization are necessary for contest-grade GPU work, but for Gated DeltaNet prefill they were not sufficient. The MSI NFER case suggests that structural reformulation, workload-shape specialization, and evaluator-aligned measurement must be treated as first-class parts of the optimization loop. This case study makes five contributions. 1. It records the final Agent-Assisted Gated DeltaNet submission by MSI NFER (MSInfer Team, 2026a;b;c),
including official speedup, approximate decode and prefill latencies, correctness coverage, and the artifact boundary used for interpretation. 2. It analyzes the team’s kernel and runtime choices, including algebraic simplification, V-dimension splitting, destination passing, cp.async, NCU-driven reasoning, and B200-oriented runtime pinning. 3. It extracts lessons from public high-performing AgentAssisted artifacts, separating decode as a shapedispatch and launch/runtime specialization problem from prefill as a structural chunk/WY/tensor-core reformulation problem. 4. It analyzes evaluator alignment and measurement tooling as optimization signals, including a Modal setting where MSI NFER decode appeared faster than FlashInfer, while an official B200 extra-round evaluation measured a much stronger baseline and reversed the conclusion. 5. It provides an evidence audit that separates official Agent-Assisted artifacts, public comparison artifacts, internal measurement logs, and companion Full-Agentstyle workflow artifacts, so that performance claims remain tied to the final submitted Agent-Assisted artifacts. Unless stated otherwise, public comparison values are drawn from the cited reports and were not independently reproduced; internal Modal measurements are reported separately from official results.
2
BACKGROUND
2.1
Contest Overview
The MLSys 2026 FlashInfer Contest was an AI kernelgeneration competition centered on submitting optimized GPU kernels for predefined FlashInfer benchmark definitions (FlashInfer Contest Organizers, 2026a;b). Participants submitted tagged Git repositories with per-definition configuration files, and the evaluator built and ran those submissions on the target platform. The contest distinguished a Full-Agent track, where the submitted agent must reproduce the kernel end to end, from an Agent-Assisted track, where human experts and agents may collaborate and the team submits the resulting kernel code. This paper concerns only the Agent-Assisted GDN submission. The GDN track had two separate definitions: decode and prefill. The contest FAQ states that the final GDN score is the average speedup across the decode and prefill definitions, and that per-definition speedup is measured against a simple reference implementation rather than an optimized FlashInfer baseline. This distinction matters for the rest of the
paper. A submission can report a favorable contest speedup while still being slower than a highly optimized FlashInfer baseline under a different evaluation setting, so we treat official score, absolute latency, correctness coverage, and benchmark source as separate pieces of evidence. 2.2
GDN Decode and Prefill
Gated DeltaNet exposes two different optimization regimes. Decode is latency-sensitive and often operates at microsecond scale. Small changes in launch overhead, synchronization, compiler target, and baseline implementation can change the apparent outcome. Prefill processes longer sequences and larger token sets. It therefore offers more opportunity for chunking, tensor-core use, and shape-specialized scheduling, but it also requires handling a recurrent state update. The MSI NFER project worked with evolving evaluation workloads. The decode workload set grew to 54 workloads, while prefill used 100 workloads. Because the benchmark suite, wrapper versions, and baseline conditions changed during the project, early internal results were not directly comparable to final official results. 2.3
Evaluation Surface
The target platform was NVIDIA B200/Blackwell. We use build and evaluation surfaces to refer to the compiler, runtime, binding, packaging, and measurement interfaces between a kernel and the evaluator. In this setting, sm 100a targeting, language choice, runtime compilation, TVM FFI, destination-passing style, and dependency pins can change both the generated code and the evaluator’s ability to load it. Our team therefore treated these interfaces as optimization concerns, not only as packaging details.
3
MSI NFER S UBMISSION
3.1
Official Result
Table 1 summarizes the final MSI NFER Agent-Assisted Gated DeltaNet result. The official headline is a 1.58× speedup. 3.2
Kernel Techniques
MSI NFER applied several useful optimizations. The shared algebraic simplification was to pre-scale state by the gate and compute output as out = q ⊤ (gS) + δ(q ⊤ k), instead of recomputing a full dot product against the updated state. This made the decode and prefill paths share a useful
Table 1. Final MSI NFER Agent-Assisted Gated DeltaNet result. Latencies are approximate per-workload averages. Metric
Value
Official Gated DeltaNet speedup Decode latency Prefill latency Decode correctness Prefill correctness Decode artifact Prefill artifact
1.58× 9.315 µs 239.48 µs 54/54 workloads 100/100 workloads 69f517b 864f286
form: g = exp(− exp(Alog ) · softplus(a + b∆t )), β = σ(b), oldv = k ⊤ (gS), δ = β(v − oldv ), S ′ = gS + kδ ⊤ . For decode, this reduced redundant state reads and unnecessary scalar synchronization. For prefill, the same form removed a second full output dot per token in the sequential recurrence. MSI NFER also used V-dimension splitting, warp-local reductions, cp.async.cg style movement for prefill K/Q staging, destination-passing style, and runtime pinning around B200 compilation. These changes improved performance in internal measurements and produced an auditable record of the optimization process. However, they did not eliminate the remaining performance gap, especially in prefill. MSI NFER did make algebraic changes; it was not a purely mechanical implementation of the baseline recurrence. However, these changes were local reformulations inside a token-by-token execution model. The public-artifact comparison in Section 5 points to a different level of computation: grouping multiple prefill tokens into chunks, summarizing the recurrence over each chunk, and exposing larger matrix-like operations that can use tensor cores. In short, MSI NFER optimized the sequential recurrence path, while the missing direction changed the computational shape of prefill. The decode basis for this report is the final AgentAssisted decode artifact at commit 69f517b, tagged gdn-decode (MSInfer Team, 2026a). We keep the artifact identity explicit because non-final and Full-Agent branches also existed, but the main text focuses on the performance interpretation rather than listing every configuration field. For the Agent-Assisted decode artifact, the team followed a
disciplined local-search process rather than adopting a late Full-Agent hybrid branch. The workflow imposed phase targets, one-change-at-a-time edits, immediate rollback on correctness or latency regression, and median-of-five promotion once Modal noise became visible. The inspected NCU log for the Agent-Assisted decode path recorded a 54/54 Modal pass and an internal arithmetic mean around 12 µs before later finalization. An NCU capture at B = 64 showed low issue-slot utilization, low achieved occupancy, and too few waves per SM; the bottleneck was insufficient parallelism and bytes in flight, not a single saturated HBM pipeline. This internal measurement is distinct from the official decode latency of 9.315 µs reported in this paper. The prefill basis is the final Agent-Assisted artifact at commit 864f286, tagged gdn-prefill (MSInfer Team, 2026b). Its evaluator-facing configuration remains relevant to reproducibility, but the main comparison concerns the type of recurrence computation that the artifact did and did not implement. For prefill, the plateau is clearer. The inspected internal and companion logs show large early gains from register-tiled state and loop fusion, from 34,280 µs to 11,510 µs, then V-split blocks from 8,260 µs to 5,210 µs, and warp-parallel V-row algebraic fusion from 5,210 µs to 2,110 µs. Later entries cluster around smaller or reverted changes: software prefetching, q/k reduction decoupling, fused qk dot, wider blocks, and additional split-factor tiers did not produce robust latency gains. This supports the main interpretation of the case study: local recurrence optimization produced real progress, then plateaued before a larger algorithmic reformulation. 3.3
What MSInfer Did Not Do
The most important limitation is that MSI NFER mostly optimized a register-resident scalar recurrence. It improved arithmetic, data movement, synchronization, and work partitioning inside that recurrence, but it did not turn prefill into a chunk-level tensor-core problem. The later public comparison shows that high-performing artifacts did. This distinction is the main reason to frame the report as a differential case study rather than a simple contest narrative. It also avoids an overclaim: MSI NFER’s algebraic simplification was meaningful, but it was a local per-token reformulation rather than the structural chunk/WY reformulation that changed the prefill workload into larger block computations.
Table 2. Decode measurement divergence. Values are approximate and should be read as an evaluator-alignment lesson, not as a controlled causal experiment. Setting
Baseline
MSI NFER
Ratio
Modal Official B200 extra-round
29.3 µs 6.55 µs
14.8 µs 8.12 µs
1.98× 0.80×
4
M EASUREMENT AND E VALUATOR A LIGNMENT
4.1
Why Modal Speedup Was Misleading
Speedup was not stable enough to be the only engineering metric. During development, a Modal run suggested that MSI NFER decode beat FlashInfer: the baseline was approximately 29.3 µs and MSI NFER was approximately 14.8 µs, or about 1.98×. An official B200 extra-round evaluation told a different story: the FlashInfer baseline was approximately 6.55 µs and MSI NFER was approximately 8.12 µs, or about 0.80× versus baseline. These numbers should not be read as a simple MSI NFER regression. They show that the reference baseline, evaluator configuration, build path, clock conditions, and hardware environment can move the denominator dramatically. The safer engineering practice is to report absolute latency, pass counts, workload coverage, build configuration, and benchmark source for every number. 4.2
Evaluation Tooling as Competitive Advantage
The measurement gap did more than complicate post hoc analysis; it changed which optimizations looked promising. MSI NFER improved its methodology with repeated runs, isolated runners, full-workload validation, and NCU profiles, but it still relied substantially on contest-facing or Modal-facing timing paths during search. Kachua’s public repository exposes a more sharply separated timing stack: run modal.py for packed benchmark/correctness runs, bench fi timing.py and bench fi timing modal.py for FlashInfer/CUPTIstyle pure GPU timing, and profile modal.py for profile artifacts. Its experiment instructions explicitly log both benchmark latency and kernel latency in microseconds and prefer the latter. This suggests that evaluation tooling itself contributed to contest performance. This distinction matters because an optimizer follows the metric it can observe. If the local harness is noisy or misaligned, a human-agent loop can select variants that win the local measurement but fail to transfer. If the harness is stable and official-aligned, smaller workload-specific choices become actionable. The MSI NFER tooling was broad but more fragmented: companion workflow tooling logged Modal
B200 latency, speedup, correctness, and workload counts, while later agent repositories added precise/decision-gate benchmark modes, NCU scorecards, surrogate scorecards, workload inventories, and phase timers. These tools helped reject bad ideas and diagnose bottlenecks, but they did not give the search loop one central FlashInfer/CUPTI-style kernel-latency primitive comparable to Kachua’s public timing path. This is the most important tooling distinction in the case study. MSI NFER did not lack engineering instrumentation; it had many forms of instrumentation. What it lacked was a single, trusted decision metric that made per-workload algorithm dispatch easy to rank. Kachua’s repository suggests a different approach: correctness and benchmark scripts still existed, but the experimental loop could privilege raw kernel timing when deciding whether a candidate formulation should advance. 4.3
Profiler Evidence
NCU profiling pushed the team away from single-number reasoning. For decode, one B200 profile showed low issueslot utilization and insufficient parallelism/bytes in flight rather than a single saturated HBM bottleneck. That matters because naive state-traffic reasoning would have labeled the kernel memory-bound too quickly. The profiler suggested a subtler small-kernel problem: too few waves and too little concurrent work for B200, especially when launch overhead and cache behavior were folded into the measurement. For prefill, profiling and comparative evidence point to a different conclusion. The issue was not only that MSI NFER needed another cache hint or fewer barriers. The missing piece was a formulation that exposes larger tensor-corefriendly chunks while preserving correctness.
5
C OMPARATIVE E VIDENCE AND L ESSONS
We now compare MSI NFER with public high-performing Agent-Assisted artifacts to extract design lessons rather than reconstruct the leaderboard. Table 3 summarizes the public top-three comparison and separates reported numbers from design lessons. We tie all exact values to pinned public reports or repositories. 5.1
Kachua: Prefill as Algorithmic Reformulation
Kachua is the most important comparison point for MSI N FER (Kachua Team, 2026a;b). Its public repository exposes multiple prefill families: Direct v2, Split-WY, and Flat-WY. The top-level kernel comment describes prefill as chunkwise matmul plus a Neumann inverse: each sequence is split into fixed-size chunks, the intra-chunk recurrence is summarized algebraically, and Tensor Core tl.dot is used for TF32/BF16/FP16 contractions. The key dispatch axes
Table 3. Public top-three Agent-Assisted GDN comparison. Reported values are drawn from the cited public artifacts. Submission
Rank
Decode
Prefill
Main design lesson
Kachua
1st
public repo raw rows: 61.17–61.50 µs
Decode: direct TVM FFI, one-warp BV=8 CUDA, vectorized loads, and butterfly reductions. Prefill: chunkwise matmul plus Neumann/WY-style inverse, adaptive CHUNK/BV, tensor-core tl.dot, and GPU-timing-aligned evaluation.
UW SyFI
2nd
public repo raw trajectory: 5.44–5.47 µs median-of-workload medians for selected CUDA variants Not reported in the inspected public artifacts
Not reported in the inspected public artifacts
LLM-CUDA
3rd
reported 6.201 µs
reported 51.992 µs
MSI NFER
outside top three
9.315 µs
239.48 µs
Decode: NVRTC/Driver API launch interface, TVM FFI, and one-warp V TILE=8. Prefill: BF16 WY-form chunked CUDA with WMMA/cp.async. The runtime interface was part of the optimization. Decode: hybrid dispatch across small recurrent, B16 high-CTA CuTe, B48 TILE V=16, and default large-batch CuTe paths. Prefill: Blackwell chunk kernel plus narrow short-shape fallback. Passed all correctness workloads. Decode stayed mostly within one templated CUDA family selected by split factor; prefill stayed in a register-resident scalar recurrence family with V-split, warp rows, gate pipeline, and local micro-optimization.
are CHUNK, the number of tokens grouped into a local recurrence solve, and BV, the number of state or output rows updated by one program. The code chooses these values from num seqs, total tokens, and an average-sequencelength proxy, then routes to Split-WY, Flat-WY, or Direct v2. This is qualitatively different from only changing block size or vector load width. The repository also gives raw microsecond evidence: prefill results.csv records accepted rows at 61.50, 61.20, and 61.17 µs, while the CUDA decode README reports selected variants with 5.44–5.47 µs median-of-workload medians. We treat these as public-repository trajectory numbers, not as independently reproduced official leaderboard averages. This comparison shows that Gated DeltaNet prefill can be improved by changing the mathematical and scheduling shape of the recurrence. Kachua converted more work into reusable chunk metadata and tensor-core-friendly local solves. MSI NFER instead made a scalar recurrence faster. That distinction explains why MSI NFER could improve internally yet remain far from top prefill latency. Kachua’s repository also shows that evaluation tooling was part of the design rather than an afterthought. The repo includes bench fi timing.py, whose header describes a FlashInfer pure-GPU timer that mirrors a CUPTI path and excludes Python dispatch and benchmark-framework overhead. Its Modal wrapper, bench fi timing modal.py, states that it matches the reference repository’s Modal CUPTI path rather than the latency path used by scripts/run modal.py, installs cupti-python, and can sweep all workloads. The same repo’s program notes say that kernel latency is more important than benchmark latency or speedup because it is the actual kernel latency. Thus the public repo supports a systems interpretation: Kachua appears to have optimized both the prefill formulation and the timing signal used to compare variants. The two reinforce each other. Structural prefill variants can be workload-sensitive, so a
kernel-latency-centered harness makes it easier to decide where a chunked, WY-style, or fallback path is actually worth keeping. 5.2
UW SyFI: Runtime Surface as a First-Class Variable
UW SyFI’s public artifacts emphasize another systems axis: the runtime and ABI surface (UW SyFI Team, 2026a;b). Their decode path uses NVRTC and the CUDA Driver API to avoid a runtime/driver mismatch, while their GDN configurations use TVM FFI and destination-passing style. Their prefill CUDA source is explicitly a BF16 tensor-core chunked kernel in WY form, with V TILE=16, C=16, four-warp block specialization, WMMA fragments, and cp.async.cg double-buffered Q/K/V loads. The inspected public artifacts do not provide a pinned table of raw decode or prefill latencies; therefore, we do not report those values. For MSI NFER, the takeaway is not simply “write CUDA instead of Python” or “use WMMA.” The deeper lesson is that the evaluator-facing surface—language, binding, driver/runtime behavior, compiler target, and output allocation convention—is part of the optimization problem. This supports MSI NFER’s own runtime-pin narrative around B200 and sm 100a, but SyFI shows that such control must be paired with a strong prefill formulation. 5.3
LLM-CUDA: Quantifying the Gap
LLM-CUDA provides the clearest public numeric comparison (LLM-CUDA Team, 2026a;b). Its retained AgentAssisted artifacts report about 6.201 µs for decode and 51.992 µs for prefill. Its prefill summary states that the retained candidate replaced an older wide pair-dispatch scheme with an upstream PR #3001 Blackwell chunk kernel as the default path, kept only a narrow recovered short-shape fallback for measured tiny-shape regressions, and passed
a 100-workload official-aligned sweep. Relative to MSI N FER ’s official numbers, this makes MSI NFER about 1.50× slower on decode and about 4.61× slower on prefill. The ratio is the main point. Decode was not solved, but prefill dominated the gap. This comparison helps avoid a misleading conclusion. If we only inspect MSI NFER’s decode story, the paper can sound like a near miss caused by evaluator noise. The prefill ratio says otherwise: the top-three path required a larger algorithmic move. 5.4
Decode Gap: Dispatch and Shape Split
The decode gap is better explained as a dispatch and smallkernel specialization gap than as a missing recurrence identity. MSI NFER’s accepted Agent-Assisted decode source is a single evaluator-facing CUDA/Python entry that compiles kernel.cu through tvm ffi.cpp.load. Inside the CUDA wrapper, batch size selects a split factor: small and mid batches use split factor=8, while larger batches use split factor=4; this maps to ROWS PER WARP instantiations of the same templated kernel. The kernel itself uses 128-thread blocks, four warps, vectorized state and q/k loads, shared staging for v, warp reductions, and L2 persistence hints (MSInfer Team, 2026a). This was a real decode implementation, but its shape split stayed within one kernel family. The public top-three decode artifacts show a more aggressive separation of regimes. Kachua’s CUDA decode README describes the retained direction as a direct TVM FFI C binding, a one-warp BV=8 kernel, 128 CTAs for B=1, vectorized state and q/k/v loads, paired butterfly warp reductions, fused decay and delta-rule logic, interleaved output/state stores, hard-coded dimensions, and launch bounds (32,1) (Kachua Team, 2026b). It also states that TVM FFI C binding was the largest speedup, that shared-memory staging hurt, and that the kernel was instruction-latency-limited rather than HBM-bandwidthlimited. UW SyFI followed a related decode shape: TVM FFI and destination passing, but with NVRTC plus the CUDA Driver API to avoid runtime-driver mismatch; the embedded decode kernel uses V TILE=8, one warp, float4 state movement, and warp-shuffle K reductions (UW SyFI Team, 2026b). These two repositories suggest that top decode work treated the launch/runtime interface and onewarp small-kernel shape as first-class variables. LLM-CUDA exposes the clearest explicit batch dispatch. Its retained decode entry point is kernel.py::kernel hybrid dispatch. The public summary routes batch size <= 8 to a Triton recurrent path, with the B8 case using a one-warp, threestage launch; batch size == 16 goes to a vendored FlashInfer CuTe pretranspose path with a higher CTA count;
batch size == 48 goes to a dedicated vendored CuTe path with TILE V=16 and four blocks per state; and the remaining large batches use the default vendored CuTe pretranspose path (LLM-CUDA Team, 2026b). The same summary reports only small but stable improvements over its same-day baseline in several regimes, which is itself informative: decode was a regime-by-regime selection problem, not a single global kernel change. This comparison reframes what MSI NFER did not do. It did try split-factor changes, 256-thread large-batch variants, shared staging, CUDA graph replay, async prefetch, cluster q/k sharing, and cache-policy variants, and many were rolled back after full-workload regressions. However, it did not converge to a top-level dispatch table with separate small, B16, B48, and large-batch paths, nor to a pure one-warp TVM-FFI/Driver-API decode design like those in the top public repositories. Its NCU evidence of low issue-slot utilization, low achieved occupancy, and too few waves per SM is consistent with this explanation: B200 was underfilled, and the winning decode direction was not only “do fewer arithmetic operations” but “choose the right launch/runtime/kernel shape for each batch regime.” This explains the 1.50× decode gap to LLM-CUDA without overstating it as the main contest gap. 5.5
Prefill Gap: Structural Versus Local Recurrence
The prefill gap was not simply a smaller version of the decode gap. Decode is a tiny-latency regime where launch overhead, underfilled SMs, and baseline drift can dominate the interpretation. Prefill exposed a more algorithmic difference. Kachua, SyFI, and LLM-CUDA all show public evidence of chunk-level prefill paths (Kachua Team, 2026b; UW SyFI Team, 2026b; LLM-CUDA Team, 2026b): Kachua uses chunkwise matmul plus Neumann/WY-style algebra and adaptive Split-WY or Flat-WY dispatch; SyFI uses a BF16 tensor-core WY-form chunked CUDA kernel with WMMA and cp.async; LLM-CUDA retained a Blackwell chunk kernel as its main path with a narrow short-shape fallback. These are not merely faster versions of MSI NFER’s kernel. They change the unit of work from one token update at a time to chunk summaries and matrixshaped subproblems. MSI NFER’s prefill log points in the opposite direction (Jang, 2026a; MSInfer Team, 2026b). It improved a registerresident scalar recurrence through state tiling, loop fusion, V-dimension splitting, warp-parallel rows, algebraic fusion, launch-bounds tuning, gate pipelining, and split-factor dispatch. The same log also records a plateau: shared-memory q/k broadcast, fast math gates, higher split factors, shuffle reductions, packed stores, software prefetching, and q/k reduction decoupling either regressed or were flat. Its own research notes concluded that tensor cores, TMA, and TMEM
were poorly matched to the current scalar recurrence, and that a step-change would require a chunked WY rewrite. This is an important nuance: within the scalar recurrence family, MSI NFER had pushed hard; the missing piece was not one more cache hint, but leaving that family. This explains the numerical asymmetry. Against LLMCUDA’s public retained artifacts, MSI NFER was about 1.50× slower in decode but 4.61× slower in prefill. Against Kachua’s public prefill trajectory rows near 61.17–61.50 µs, MSI NFER’s 239.48 µs prefill was roughly 3.9× slower. We therefore treat decode as an important but secondary gap, and prefill as the main performance gap caused by missing structural reformulation, tensor-core exposure, and workload-specialized dispatch.
kernel-latency harness. 6.4
The prefill gap and the measurement gap should not be treated as independent failures. Chunked and WY-style formulations introduce choices about sequence regime, chunk size, fallback path, and tensor-core-friendly tiling. Those choices are only useful if the benchmark loop can tell which shapes improved. The top public artifacts suggest that algorithm reformulation and evaluator-aligned timing coevolved; MSI NFER instead had strong local evidence collection but a weaker central signal for promoting workloadspecific structural changes. 6.5
6
L ESSONS L EARNED
6.1
Decode Was Shape-Dispatch; Prefill Was Structural
Algorithm and Measurement Co-Design
Rejected Experiments Are Evidence
The strongest lesson is that performance in the two tasks plateaued for different reasons. Decode improvements mattered, and the public repositories suggest that the missing decode move was stronger batch-regime dispatch and launch/runtime specialization. But public ratios show that MSI NFER was much farther from top-three prefill than topthree decode. Future AI-assisted kernel systems should therefore ask two questions earlier: whether a microsecondscale decode kernel needs separate shape-specific launch paths, and whether a prefill operator needs algorithmic reformulation before many iterations are spent on local recurrence micro-optimizations.
The repo-local logs are useful precisely because they record failed hypotheses. They show that plausible local changes can be wrong for different reasons: decode cache-policy and cluster-sharing ideas could regress full-workload latency even when they looked reasonable on paper; decode NCU evidence showed that B200 was underfilled rather than purely bandwidth-saturated; prefill software prefetching added overhead despite high L1 hit rate; and q/k reduction decoupling could not overlap work that still executed in warp program order. Separate Full-Agent branches also showed a different failure mode: a candidate can be technically interesting yet still violate the evaluator’s submission contract. These rejected experiments strengthen the casestudy claim. The absence of stronger prefill performance was not simply a lack of trying small CUDA edits; the remaining gap was structural.
6.2
6.6
Build Surface Is Part of Kernel Optimization
Compiler target, language, binding, runtime/driver behavior, and destination-passing style can affect both performance and evaluator compatibility. SyFI’s NVRTC/Driver API path and MSI NFER’s sm 100a/runtime-pin concerns point to the same conclusion: the evaluator-facing build surface is a real systems variable. 6.3
Measurement Is an Optimization Surface
The Modal/official divergence shows that benchmark configuration can steer the search toward or away from useful candidates. For microsecond-scale decode, repeated medians, isolated runners, per-workload views, FlashInfer/CUPTIstyle pure GPU timing, and official-parity validation are not bookkeeping; they are part of the optimization method. In retrospect, MSI NFER’s evaluation toolchain was useful but less centralized than the strongest public timing workflow: it reduced noise and added profiler context, but it did not fully replace the template path with an official-aligned
Full-Agent-Style Companion Evidence
Full-Agent experiments were also attempted, but an error in the final submission/evaluator path prevented an official result. We therefore use those artifacts only as evaluator-contract and search-process context, not as AgentAssisted performance evidence. The associated repo-local workflow skills—research, optimize, bench, and log-result (Jang, 2026b)—document search behavior and rejected hypotheses; official Agent-Assisted performance evidence remains limited to decode 69f517b and prefill 864f286. In the inspected jsg1504/mlsys26 workspace, these skills made the optimization trace explicit. The loop read the current kernel, checked previous attempts, collected or requested profiler evidence, implemented one change, benchmarked it on Modal B200, and appended structured records. By the end of the inspected companion log, the prefill history contained 33 benchmark records, and the broader workspace also contained decode optimization records.
Table 4. Repo-local companion automation evidence loop inspected for this case study. This is not official Agent-Assisted submission evidence. Unit
Evidence produced
Reads kernel code, benchmark history, and optimization log; returns ranked ideas with expected impact and risk. optimize Runs a one-change loop: profile, select, implement, benchmark, profile again, then accept or revert. bench Packs the solution and runs Modal B200 benchmarks, recording status, latency, speedup, correctness, workload count, and notes. log-result Summarizes benchmark trends and optimization-history learnings from JSONL and Markdown logs. research
8
C ONCLUSION
Our team, MSI NFER, produced a correct Agent-Assisted Gated DeltaNet submission with an official 1.58× speedup in the MLSys 2026 FlashInfer Contest. Its value lies not in its final rank but in the boundary it exposes. AI-assisted micro-optimization, runtime pinning, and careful measurement can produce a useful kernel artifact. Yet the public top-three submissions show that for Gated DeltaNet prefill, the top path required a larger algorithmic shift: chunked, shape-specialized, tensor-core-friendly recurrence computation. The broader lesson is that this algorithmic shift also required a reliable timing signal to select among workloadspecific variants. In this sense, the kernel body was necessary but not sufficient.
R EFERENCES This loop made the search auditable, but its primary benchmark records were still Modal benchmark latency and speedup records. Later MSI NFER agent repositories added more specialized analysis—NCU scorecards, surrogate scorecards, workload inventories, and direct phase timers— which improved diagnosis. The evidence base was therefore broad, but not organized around a single official-adjacent kernel-latency objective. That distinction helps explain why many local micro-optimizations were explored, while the larger workload-specific prefill reformulation path remained harder to evaluate and promote. 6.7
AI Assistance Still Needs Systems Judgment
Agents can generate code, summarize logs, and propose variants. They do not automatically know which measurements transfer to the official evaluator, whether a prefill recurrence needs a WY reformulation, or whether a package satisfies the evaluator’s source-discovery contract. Human judgment remained central to deciding what evidence counted.
7
L IMITATIONS
This report is a case study, not a reproduced benchmark of all public submissions. We tie public comparison values to repository commits or official writeups. Several interpretations of root causes are based on logs and public comparison rather than isolated controlled experiments. We therefore use careful language: “suggests,” “indicates,” and “is consistent with” are often more accurate than causal absolutes. This report intentionally keeps the main story focused on the Agent-Assisted GDN submission. Full-Agent experiments belonged to a separate track and are not used as performance evidence for the Agent-Assisted result. Where they are mentioned, we frame them only as cautionary examples about submission and evaluator contracts.
FlashInfer Contest Organizers. MLSys 2026 FlashInfer contest materials and leaderboard. https://github. com/flashinfer-ai/mlsys26-contest, 2026a. Official contest site and leaderboard; accessed 2026-07-18. FlashInfer Contest Organizers. MLSys 2026 FlashInfer contest starter kit and faq. https://github.com/flashinfer-ai/ flashinfer-bench-starter-kit, 2026b. FAQ and evaluation instructions; accessed 2026-07-18. Jang, J. MLSys 2026 companion GDN optimization logs. https://github.com/jsg1504/mlsys26/ tree/optimize-kernel/logs, 2026a. GitHub repository under jsg1504; companion prefill benchmark history and optimization log inspected; not official Agent-Assisted submission evidence; accessed 2026-06-07. Jang, J. MLSys 2026 optimization workflow skills. https://github.com/jsg1504/mlsys26/ tree/optimize-kernel/.claude/skills, 2026b. GitHub repository under jsg1504; workflow units inspected: research, optimize, bench, and log-result; accessed 2026-06-07. Kachua Team. Gated DeltaNet agent-assisted writeup. https://github.com/flashinfer-ai/ mlsys26-contest/blob/main/writeups/ 2026a. Kachua-GDN-Agent-Assisted.pdf, MLSys 2026 FlashInfer Contest artifact; accessed 2026-06-07. Kachua Team. Kachua MLSys 2026 FlashInfer contest repository. https://github.com/romitjain/ kachua-mlsys, 2026b. Accessed 2026-06-07.
LLM-CUDA Team. LLM-CUDA agent-assisted GDN writeup. https://github.com/ flashinfer-ai/mlsys26-contest/blob/ main/writeups/LLM-CUDA_AgentAssisted. pdf, 2026a. MLSys 2026 FlashInfer Contest artifact; accessed 2026-06-07. LLM-CUDA Team. MLSys 2026 FlashInfer contest repository. https://github.com/syhya/ mlsys26-flashinfer-contest, 2026b. Accessed 2026-06-07. MSInfer Team. MSInfer agent-assisted GDN decode artifact. GitHub commit 69f517b, 2026a. Agent-Assisted decode artifact at 69f517b; accessed 2026-07-18. MSInfer Team. MSInfer agent-assisted GDN prefill artifact. GitHub commit 864f286, 2026b. Agent-Assisted prefill artifact at 864f286; accessed 2026-07-18. MSInfer Team. MSInfer GDN agent-assisted decode and prefill repository. https://github.com/ Bammuri/mlsys26, 2026c. Accessed 2026-06-07. UW SyFI Team. UW SyFI agent-assisted GDN writeup. https://github.com/flashinfer-ai/ mlsys26-contest/blob/main/writeups/ UW_SyFI_AgentAssisted.pdf, 2026a. MLSys 2026 FlashInfer Contest artifact; accessed 2026-06-07. UW SyFI Team. UW SyFI agent-assisted MLSys contest repository. https://github.com/kamahori/ mlsys-contest-syfi-agent-assisted, 2026b. Accessed 2026-06-07.
A
E VIDENCE AUDIT
Table 5. Evidence sources and status for reported claims. Claim
Evidence class
The final Agent-Assisted Gated DeltaNet submission by Official result MSI NFER achieved a 1.58× speedup Contest material Contest GDN scoring averaged decode and prefill definition speedups, and speedup was measured against the simple definition reference rather than an optimized FlashInfer baseline Final decode latency 9.315 µs and prefill latency 239.48 µs Official result Decode passed 54/54 and prefill passed 100/100 workloads LLM-CUDA reported 6.201 µs decode and 51.992 µs prefill
Official / artifact Public report
Kachua public repository records selected decode variants at 5.44–5.47 µs median-of-workload medians and prefill rows at 61.17–61.50 µs Kachua decode source and README describe direct TVM FFI C binding, one-warp BV=8 CUDA, vectorized loads, paired butterfly reductions, and launch bounds (32,1) Kachua prefill source uses chunkwise matmul, Neumann/WY-style algebra, adaptive CHUNK/BV, Split-WY/Flat-WY dispatch, and tensor-core tl.dot contractions SyFI used NVRTC/Driver API, TVM FFI, DPS, a one-warp V TILE=8 decode kernel, and a BF16 tensor-core WY-form chunked prefill kernel with WMMA/cp.async LLM-CUDA retained decode used hybrid dispatch: B <= 8 recurrent path, B16 high-CTA CuTe path, B48 TILE V=16 path, and default large-batch CuTe path LLM-CUDA retained prefill replaced wide pair dispatch with a Blackwell chunk kernel as the main path and a narrow short-shape fallback Modal vs official decode divergence values
Public repo
Kachua repository contains bench fi timing.py, bench fi timing modal.py, and program notes that distinguish kernel latency from benchmark latency and prefer the former MSI NFER agent repositories and companion JSG workflow tooling logged Modal benchmark latency/speedup histories and later added NCU, surrogate-scorecard, workload-inventory, and phase-timer analysis; this audit did not find a single FlashInfer/CUPTI-style kernel-latency objective used as the central optimization primitive Algorithm reformulation and evaluator-aligned timing appear to have reinforced each other in public high-performing prefill workflows Repo-local companion automation loop used research, optimize, bench, and log-result workflow units; this loop is not official Agent-Assisted submission evidence Decode artifact is the Agent-Assisted gdn-decode state at 69f517b Decode config uses language = "cuda", decode submit entry.py::run, source files kernel.cu and decode submit entry.py, and destination-passing style MSI NFER accepted decode selected split factor by batch size inside one templated CUDA kernel family rather than a public high-performing hybrid dispatch table Prefill artifact is the Agent-Assisted gdn-prefill state at 864f286 Prefill config uses language = "python", msinfer entry.py::run, and destination-passing style for gdn prefill qk4 v8 d128 k last Full-Agent experiments were attempted, but an error in the final submission/evaluator path prevented an official result; those artifacts are excluded from Agent-Assisted performance evidence Companion prefill benchmark history contained 33 records
Companion prefill logs show early large local-recurrence gains followed by reverted or flat micro-optimizations and an explicit scalar-recurrence plateau Prefill reformulation was the largest visible gap
Public repo
Evidence status or source Final leaderboard or evaluator output available to the authors; reported here as official, not independently reproduced. Contest FAQ and starter-kit material (FlashInfer Contest Organizers, 2026b); the interpretation is scoped to that contest documentation. Per-definition final evaluator output available to the authors; reported as approximate averages. Final validation logs for the submitted decode and prefill artifacts. LLM-CUDA writeup and retained-artifact repository (LLM-CUDA Team, 2026a;b); not reproduced here. Kachua decode README and prefill results.csv (Kachua Team, 2026b); treated as public trajectory numbers, not official leaderboard averages. Kachua decode README and kernel.cu (Kachua Team, 2026b).
Public repo
Kachua prefill kernel.py, results.csv, and experiment log (Kachua Team, 2026b).
Public repo/report
SyFI source, configs, and writeup (UW SyFI Team, 2026a;b); no raw latency table is reported here.
Public repo/report
LLM-CUDA retained decode summary and source (LLM-CUDA Team, 2026a;b).
Public repo/report
LLM-CUDA retained prefill summary and source (LLM-CUDA Team, 2026a;b).
Internal / official
Internal Modal logs and official extra-round evidence available to the authors; used as an evaluator-alignment example, not a controlled causal experiment. Kachua timing scripts and repository notes (Kachua Team, 2026b).
Public repo
Artifact / interpretation
Inspected jsg1504/mlsys26 skills/logs, MSI NFER agent scripts, benchmark logs, and profiler artifacts; negative claim is limited to the inspected artifacts.
Interpretation
Derived from public repository comparison and timing-script evidence; not a controlled ablation or same-harness reproduction.
Artifact / workflow context
jsg1504/mlsys26 workflow skills and logs (Jang, 2026b;a); companion context only.
Artifact
MSInfer Agent-Assisted decode repository metadata (MSInfer Team, 2026a). Checked against the local Agent-Assisted decode config.toml and source tree.
Artifact
Artifact / interpretation
Agent-Assisted decode kernel.cu, rejected-experiment logs, and public decode-source comparison; not a claim of exhaustive search.
Artifact
MSInfer Agent-Assisted prefill repository metadata (MSInfer Team, 2026b). Checked against the local Agent-Assisted prefill config.toml and source tree.
Artifact
Internal / artifact
Internal validation or submission records; included only as evaluator-contract context, not Agent-Assisted performance evidence.
Artifact / workflow context
Verified from logs/prefill/bench history.jsonl; companion evidence, not official Agent-Assisted submission evidence. Prefill optimization log and benchmark-history entries (Jang, 2026a); companion evidence only.
Internal measurement / interpretation Interpretation
Supported by public ratios and design comparison; not a same-harness reproduction of all submissions.