# How two Nodus engineers and 16 ByteAsk agents raised decode throughput 2.3x and training throughput 52% in 30 days

The first pull request moved the benchmark's solo reference from decode's own 33 SMs to all 132, which put co-located decode at 33% of solo instead of 61%. Across 2,140 hardware-gated experiments on the same H100s, Nodus merged 226 pull requests, and the projected fleet for these workloads fell from six 8-GPU nodes to four.

Company

Nodus Compute, a multi-tenant GPU platform for training, RL and inference

Workload

Qwen3-4B shapes in BF16; 16 concurrent LoRA rank-8 training jobs on the q and v projections; decode requests of 1 to 8 sequences with contexts of 512 to 8,192 tokens, continuously batched

ByteAsk

16 agents (ByteAsk Harness with Context, Claude Opus 5.5) for 30 days in Nodus's agent sandboxes

Hardware

One 8 x H100 SXM 80 GB node for measurement, one exclusive GPU per gated run

Human time

Two Nodus engineers: 150 hours of review, 40 hours of setup and supervision

Nodus runs LoRA training and decode serving inside one persistent CUDA kernel on an H100. For 30 days, 16 ByteAsk agents ran continuously against that stack, and two Nodus engineers reviewed their output.

Nodus's benchmark compared co-located decode with solo decode on the 33 SMs decode has to itself, and reported 61%. The first pull request Nodus merged from ByteAsk moved the solo reference to all 132 of the H100's SMs, against which co-located decode stood at 33%. Over the next 225 merged pull requests, co-located decode rose from 1,180 to 2,710 tokens/s at a 50 ms p99 inter-token latency, 76% of that reference. Co-located training throughput rose 52% over the same period, and training's efficiency relative to a solo run of the same build went from 79% to 85%. Nodus's capacity plan for these workloads dropped from six 8-GPU nodes to four.

## 1. What Nodus runs

Nodus operates a GPU platform on which customer training and low-latency serving share the same H100s. Each GPU is divided in two ways.

Inside the persistent kernel, decode has 33 dedicated SMs and training has the other 99. When decode requests queue up, as many as 33 training SMs finish their current tile, switch to decode work and then return to training. Decode outputs have to match a solo decode run bit for bit, so tile shapes and split-KV factors depend only on batch size and sequence length, never on how many SMs are on loan at the time.

Outside the kernel, a host-side limiter confines each tenant's CUDA streams to an assigned set of TPCs and rate-limits its kernel launches with a token bucket.

> “A bad change to this kernel rarely crashes. A tile boundary moves, a reduction runs in a different order, and one output bit changes in one layout out of a hundred. Every change has to pass the solo-versus-co-located check before anyone reads it, which is why we were careful about letting agents near it.”

Viswa Kotra, CTO, Nodus Compute

## 2. What ByteAsk added to the model

All 16 agents ran Claude Opus 5.5. The model wrote the candidate CUDA changes. ByteAsk chose the code each agent worked on and decided when a change was finished.

Context indexed the H100 build configuration, so when a kernel existed in H100 and B200 variants behind build flags, the agent worked on the variant the H100 build compiles. The repository has 431 CUDA launch sites across the persistent kernel, the reference and test kernels, the benchmarks and the limiter's tests. Context resolved 412 of them, including launches made through Nodus's launch helpers and function-pointer tables, and flagged the other 19 as unresolved.

The Harness ran every candidate on an exclusive H100 against Nodus's gates before a person saw it, and discarded the ones that failed. Reviewer comments were attached to the agent's next attempt on related code.

The Cloud Agent queue kept the 16 agents scheduled for 9,800 of the 11,520 available hours. About 5,500 of those hours went to editing, building and profiling; the rest was spent waiting for gate results.

Section 4 measures how much the Harness and Context each add to the same model.

## 3. Measuring against the whole GPU first

Nodus's dashboard reported co-located decode at 61% of solo throughput, with the solo run on the 33 SMs decode has to itself inside the persistent kernel. The first pull request added a solo reference on all 132 SMs: 3,570 tokens/s against 1,935 on 33 SMs, at a median of 28 sequences in flight. Co-located decode stood at 33% of it, and every later change was measured against it.

The same week, the agents made 22 changes that readied the gates for unattended agent runs, each with a fixture that fails without it. The scorer now checks each run's exit status and its log for CUDA errors, which a reviewer had done by reading the log.

The agents also built a nightly benchmark suite of 160 workloads covering request size, context length and tenant mix. It runs on one node in 38 minutes. Nodus engineers ran their own pull requests through it as well, and during the month it stopped 9 of them before merge, each of which would have cut throughput or raised p99 latency by more than 3%.

> “The first pull request we merged took our dashboard from 61% to 33%. It was also the reason we trusted the next 225.”

Pranav Varshney, CEO, Nodus Compute

## 4. The same model, with and without ByteAsk

In the second week, Nodus pulled 40 tasks from the queue that no agent had started: 22 performance targets and 18 correctness checks on the code paths the study was changing. Nodus engineers later found that 15 of the checks needed a code change and 3 showed correct behavior.

Each task went to three setups. All three used an identical task statement, repository snapshot, starting commit, Claude Opus 5.5 configuration, shell, compiler, profilers and test infrastructure, with a per-task budget of 6 million model tokens and 4 H100-hours, and made three independent attempts per task. Both ByteAsk setups included the Harness, and the full setup also enabled Context. Each setup was isolated from the others' attempts, review comments and later repository states, and these tasks were excluded from the main agent queue until the comparison finished.

|  | Claude Code | ByteAsk Harness | ByteAsk Harness + Context |
| --- | --- | --- | --- |
| Tasks correctly resolved, of 40 | 13 | 21 | **32** |
| Actionable tasks with a change that passed every gate, of 37 | 13 | 19 | **29** |
| Checks closed with evidence that no change was needed, of 3 | 0 | 2 | **3** |
| False defect reports, of 120 attempts | 11 | 7 | **3** |
| Attempts that edited code the H100 build does not compile, of 120 | 17 | 9 | **0** |
| Median time to the first gated change, on the 13 tasks all three solved | 3.4 h | 2.1 h | **0.9 h** |

A task counts as resolved if the setup produced a change that passed every gate or, when a check needed no change, a reproducible test showing the code behaved correctly. The Harness reduced false defect reports from 11 to 7 because it does not report a defect until a gate reproduces it. Context reduced them to 3. Context also eliminated edits to code that the H100 build does not compile, with none across 120 attempts compared with 17 for Claude Code and 9 for the Harness alone.

## 5. The 30-day loop

Each night the agents profiled the kernel and the serving path with Nsight Compute and Nsight Systems, turned what they found into candidate changes, and ranked the candidates by expected gain per GPU-hour. The Harness placed every candidate on an isolated branch with a test and ran it through a 6-minute smoke gate of 12 workloads on one H100. A candidate that passed ran the full 160-workload suite before it became a pull request.

| Stage | Count |
| --- | --- |
| Experiments | 2,140 |
| Failed a correctness gate | 1,020 |
| Performance candidates that passed correctness but gained less than the 3% minimum effect or stayed inside measurement noise | 430 |
| Passed every gate | 690 |
| Folded into a larger change or superseded by a better one | 278 |
| Pull requests | 412 |
| Merged | 226 |
| Declined: overlapped a better change (71), conflicted with a planned redesign (58), gain too small for the added complexity (39), maintainability (18) | 186 |

Of the 226 merged pull requests, 64 were kernel performance changes, 51 extended the limiter or hardened synchronization, 58 added tests and benchmark workloads, 29 changed serving and scheduling, and 24 changed build and tooling. The median pull request changed 84 lines (410 at the 90th percentile) and took 18 minutes of active review. Together they touched 214 files in 9 subsystems. Reviewers re-ran the measurements themselves on 1 in 5 pull requests and on every change to the limiter or to synchronization.

## 6. Results

| Change | Measured effect |
| --- | --- |
| Decode matrix-vector kernels specialized per batch size | per-layer MLP down-projection time at batch 1: 61 µs to 18 µs |
| Split-KV decode attention with a fixed two-stage reduction | p99 inter-token latency at 8k context: 41 ms to 24 ms |
| Lending that claims a training tile only when its dependencies are ready | decode p99 under load: -33%; co-located training stalls: -7% |
| Training GEMMs moved from mma.sync to Nodus's WGMMA and TMA backend, with register reallocation between producer and consumer warpgroups (setmaxnreg) for a 4-stage pipeline | training GEMM throughput: 1.5x |
| Staged, coalesced loads of LoRA down-projection inputs | LoRA adapter-gradient kernel: -29% time; full training step: -4% |

With all changes in, co-located decode at the 50 ms p99 inter-token limit rose from 1,180 to 2,710 tokens/s. Over 30 paired measurements, the bootstrap 95% interval for the 2.3x gain was 2.2x to 2.4x. That is 76% of what the original kernel delivered with the whole GPU and no training, up from 33%. The changes also sped up solo decode, from 3,570 to 4,820 tokens/s, so measured against the final build running solo, co-located decode went from 33% to 56%.

Co-located training throughput rose 52%. Most of that came from the training GEMMs. Nodus had already written a WGMMA backend that passed its cuBLAS check, but the training path still used mma.sync, and the agents moved it over. Solo training improved 42% from the same changes, so co-located training went from 79% to 85% of a same-build solo run.

Decode logits stayed bitwise identical between solo and co-located runs on all 160 workloads under greedy BF16 decoding, and training stayed within Nodus's PyTorch reference tolerances.

### One change in detail

At batch 1, the decode matrix-vector kernel allocated eight accumulators per thread and ran shuffle reductions for all eight, regardless of batch size. On the MLP down-projection, which reads a 49.8 MB BF16 weight per layer, Nsight Compute showed 168 registers per thread, 25% occupancy and DRAM throughput at 24% of peak. The agent templated the kernel on batch size without changing the reduction order or weight reuse, so outputs stayed bitwise identical. Registers dropped to 72, occupancy rose to 50% and DRAM throughput reached 83% of peak. Time per layer fell from 61 to 18 µs. The patch, both profiles and the command that reproduces the result on an H100 are part of the data release.

## 7. The limiter

Fifty-one of the merged pull requests extended the host-side limiter or hardened the kernel's synchronization. Among them:

- The agents extended launch-rate enforcement to kernels replayed through CUDA graphs.
- A launch waiting on the token bucket when its tenant's quota is cut resumes at the new rate.

Each change has a host-side test that fails without it. Confinement tests on an H100 launched kernels that recorded their SM IDs under quota changes, graph replays and contention, confirming that no kernel ran outside its assigned TPCs. A fuzzer written by the agents ran 10 million randomized sequences of quota changes, launches and graph replays against a host model of the limiter, checking after every step that no tenant exceeded its TPC set or launch quota. 2,000 of those sequences were replayed on an H100. These tests cover compute allocation only; memory isolation was out of scope.

## 8. Economics

| Duration | 30 days |
| --- | --- |
| Merged pull requests | 226 |
| Agent time | 9,800 scheduled agent-hours |
| Human time | 190 hours (150 review, 40 setup and supervision) |
| Model spend | $41,600 |
| Measurement GPU use | 4,300 H100-hours |
| Reserved GPU capacity | 5,760 H100-hours (one 8 x H100 node for 30 days) |
| Reserved GPU cost | about $13,800 at Nodus's internal accounting rate of $2.40 per H100-hour |
| Human cost | about $38,000 at a fully loaded $200 per engineer-hour |
| Total cost | about $93,400, or $413 per merged pull request |

Nodus's November capacity plan uses its production request mix and latency targets. With the measured changes, the capacity plan for these co-located workloads requires 32 H100s instead of 48, reducing the allocation from six 8-GPU nodes to four. Decode alone would fit on 21 GPUs, so committed training throughput determines the fleet size. At Nodus's internal rate of $2.40 per H100-hour, 16 fewer planned H100s represent about $28,000 of capacity a month, or $336,000 a year, against a total study cost of about $93,400. This is a planning result. No deployed GPUs were removed during the study.

## 9. Limits

The results come from one codebase, one model shape and one GPU type over 30 days. B200 paths were out of scope. Nodus's engineers set the architecture and the optimization priorities, and they declined 186 changes that had passed every gate.

> “Two of us reviewed everything that merged, and review was the slow part, not the agents. Next we are pointing it at checkpoint and restore, where the hard part is the same: proving it stays correct under load.”

Viswa Kotra, CTO, Nodus Compute

## 10. A pilot on your stack

A ByteAsk pilot runs on one branch and one GPU node and begins with an audit of the team's correctness and latency gates in the first week. The team receives every patch, its measurement logs and all rejected experiments, and decides what to merge.

[Book a demo](https://byteask.cal.com/pratyush/book-a-demo) [How ByteAsk handles your code](https://byteask.ai/enterprise#security)

## Appendix A. Method

- **Headline workload.** Derived from Nodus's October production serving trace: each arrival is a request of 1, 2, 4 or 8 sequences with probabilities 40%, 30%, 20% and 10%; context lengths log-uniform from 512 to 8,192 tokens; 256 output tokens per sequence; Poisson arrivals; continuous batching on the server. At the 50 ms limit the median number of sequences in flight was 28 for solo decode and 21 for co-located decode after the run. Training runs continuously beside decode: 16 LoRA rank-8 jobs on the q and v projections of Qwen3-4B shapes, sequence length 2,048.
- **Throughput at the latency limit.** For each build, offered load was increased in 5% steps until the p99 inter-token latency exceeded 50 ms. Inter-token latency is the time between consecutive output tokens of a sequence; the first token (time to first token) is excluded, and the p99 is computed over every interval in a 10-minute window. The reported throughput is the highest load that stayed under the limit.
- **Pre-run solo decode reference.** The original build ran the same decode workload on all 132 SMs with no training present: 3,570 tokens/s.
- **Final-build solo decode reference.** The final build under the same conditions: 4,820 tokens/s.
- **Solo training reference.** The same 16 LoRA jobs on all 132 SMs with no decode present, measured on the same build as the co-located run it is compared with.
- **Agents.** ByteAsk CLI 0.2.4 with the CUDA-capable Context engine, Claude Opus 5.5 at medium effort, in Nodus's agent sandboxes. Comparison setups as in section 4: Claude Code 2.1.322, ByteAsk Harness without Context, and ByteAsk Harness with Context, same model and effort, each with its own budget per task. No network beyond the model API in any setup.
- **Experiment.** One candidate change on an isolated branch with its test, a smoke gate on one exclusive H100 and, if it passed, the full suite. Counted from first agent action to gate verdict.
- **Gates.** Bitwise equality of decode logits between solo and co-located execution under greedy BF16 decoding; training loss, parameters and gradients against a PyTorch reference (relative tolerance 1e-3 for loss, 1e-2 for parameters and gradients); for performance candidates, throughput or p99 on the 160-workload suite with a 3% minimum effect threshold; for correctness candidates, a reproducer that fails before the change and passes after it; TPC confinement tests for limiter changes; additional shapes not visible to the agents for kernel changes. Every candidate also passed the applicable correctness gates.
- **Measurement.** H100 SXM 80 GB, CUDA 12.8, driver 575.57, SM clocks locked at 1,755 MHz. Warm-up, then 30 paired measurements per reported change over three sessions in randomized order. Bootstrap 95% intervals: co-located decode at the 50 ms limit 2.3x (2.2x to 2.4x); co-located training throughput +52% (+48% to +56%); batch-1 down-projection 3.4x (3.2x to 3.6x).
- **GPU-hour budget.** Smoke gates 214 H100-hours; full-suite runs before each pull request 2,089; nightly suites 152; headline measurement campaigns 1,500; profiling and reviewer re-runs 345; total 4,300.
- **Capacity model.** Nodus's November plan sized co-located nodes to serve its forecast decode load at a 50 ms p99 inter-token latency while delivering its committed training throughput. With the measured per-GPU decode capacity and training throughput, decode alone requires 21 H100s and training requires 32, bringing the planned fleet to 32 GPUs across four 8-GPU nodes instead of six. Nodus ran the model; inputs and outputs are in the data release. This is a planning result; no deployed GPUs were removed during the study.
- **Checks that needed a code change.** Counted when a Nodus engineer confirmed the change was needed and a test failed before it and passed after.
- **False defect report.** A suspected defect reported by an agent that Nodus engineers could not reproduce and for which no failing gate or test case could be produced.
- **Data release.** Patches, gate logs, raw timings, build flags and profiler captures for every reported number will be released with Nodus's permission.

[All customer stories](https://byteask.ai/customers)

---
HTML version: https://byteask.ai/customers/nodus