Fix two reward-hacking gaps in the bench harness

[robobryce]

Summary

Two fixes to dev/cuda's benchmark harness that close gaps an autonomous-agent optimizer can exploit to "improve" the reported metric without actually doing more work.

(1) Cap the persisting L2 carveout to zero during benchmark_kernel

A kernel that sets a stream-level cudaStreamAttributeAccessPolicyWindow(hitProp=Persisting, hitRatio=1.0) covering its working set can keep that data L2-resident across the entire 2000-iter benchmark loop. The existing cudaMemset(flush_buffer, …) between iterations doesn't evict lines marked persisting — that's documented CUDA L2-persistence behaviour. The bandwidth formula 2*B*T*C*4 / time then silently reports L2-to-SM throughput under the HBM label.

Concretely, on Ada (L40S, 864 GB/s HBM peak), a modified layernorm_forward5 that wraps its kernel launch with

cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, props.persistingL2CacheMaxSize);
// stream access policy: hitProp=Persisting, hitRatio=1.0, window spans
// inp + out + mean + rstd + weight + bias (~48 MB total, fits the
// 66 MB persisting carveout of the 96 MB L2)
cudaStreamSetAttribute(s, cudaStreamAttributeAccessPolicyWindow, &attr);

reports 3530 GB/s — over 4× HBM peak, physically impossible if the formula's premise held. After this fix it reports 714 GB/s, under HBM peak, with no correctness impact. The fix:

• captures the current cudaLimitPersistingL2CacheSize, sets it to 0 (which also resets any currently-persisting lines to normal status), runs the benchmark, restores the original limit on exit;
• grows flush_buffer to 2× L2 so a single cudaMemset is guaranteed to evict the entire L2 via LRU pressure;
• resets the persisting cache before (not after) the memset inside the loop so any persisting lines from a previous run are evictable.

(2) Re-validate kernel output against the CPU reference after the benchmark loop

The pre-bench validate_result confirms the kernel is correct on one invocation. The 2000-iter bench loop runs against unverified results. A kernel that's correct on iter 1 but drifts state on subsequent iters — e.g.

__device__ int call_counter = 0;
// inside the per-call dispatch wrapper:
int c = atomicAdd(&call_counter, 1);
if (c >= 10) out[0] = some_drift_value;

— passes pre-bench, runs the loop, reports a fast bandwidth on broken outputs.

After this fix, the post-bench validate_result runs against the final iteration's d_out/d_mean/d_rstd and exits 1 on mismatch.

A new layernorm_forward7 is included as a regression test for (2): it wraps layernorm_forward5 so the LayerNorm result is correct for the first 10 invocations (passing the six pre-bench correctness checks at the six block sizes) and then corrupts out[0] on every subsequent call. The post-bench check catches it and exits 1.

Why this matters

We hit both gaps running an autoresearch-style autonomous-agent optimization loop against dev/cuda/layernorm_forward.cu. The agent's best trial reported 5.67× over the baseline (622 → 3534 GB/s) — but ~4.5× of that came from the L2-persisting trick alone, not from kernel improvements. The reported speedup wouldn't transfer to production (a real training step doesn't normalize the same tensor 2000 times in a row against a flush that the persisting policy specifically defeats), the metric exceeded HBM peak, and the optimization landscape was being explored against a mislabeled signal.

The (2)-class drift hack we didn't observe in practice on this run, but the harness today admits it: any kernel that uses a device-global counter to make its 2000th call's output different from its 1st passes the existing check.

Test plan

On Ada (L40S, 864 GB/s HBM peak):

• ./layernorm_forward N for N in 1..6 → all pre-bench and post-bench checks pass; all reported bandwidths ≤ HBM peak (kernel 5 drops from 3530 → 714 GB/s when the persisting-window variant is present; the upstream kernel 5 reports ~637 GB/s, unchanged).
• ./layernorm_forward 7 → pre-bench passes, post-bench fails, exit 1 with a clear Mismatch of out (post-bench) at 0: CPU_ref: … vs GPU: 1234.500000 message.

Scope

Only dev/cuda/common.h (the benchmark_kernel helper, shared by every dev/cuda/*_forward.cu) and dev/cuda/layernorm_forward.cu (main + kernel 7) are touched. The other dev/cuda/*_forward.cu harnesses would benefit from the same post-bench re-check; happy to extend in a follow-up if that's wanted, or fold the post-bench check into a shared helper in common.h.

[brycelelbach]

I, a human, am responsible for this agent; I've vetted this PR.