Harden L2 cache clearing against persistence reward-hacking

[robobryce]

Summary

clear_cache() (csrc/clear_l2.cu) evicts the previous iteration's L2 footprint by spam-writing 2 * L2 of dummy data, and calls cudaCtxResetPersistingL2Cache() to undo "sneaky cache persistency". The reset on its own is not sufficient when a submission also reserves a persisting L2 carveout. This PR closes that gap and adds a regression test.

The hole

A kernel can keep a large reused input (weights, lookup tables, embedding rows, histogram bins, ...) resident in L2 across the timed iterations by marking it persisting:

cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, N);      // reserve a set-aside L2 carveout
cudaStreamAttrValue w{};
w.accessPolicyWindow.base_ptr  = reused_buffer;
w.accessPolicyWindow.num_bytes = N;
w.accessPolicyWindow.hitRatio  = 1.0f;
w.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &w);   // on the benchmark stream

Normal / streaming writes can only use the non-set-aside portion of L2, so the write_cache_kernel spam-write cannot evict lines parked in the carveout. The reused buffer then stays warm across iterations and the kernel is timed as if its inputs were permanently L2-resident — a benchmark score that a fair cold-cache (or real one-shot) run would never see. cudaCtxResetPersistingL2Cache() is meant to neutralize this by demoting persisting lines to "normal", but it is the only thing standing between the carveout and the spam-write, and it is a host-immediate, non-stream-ordered call.

Two aggravating facts I confirmed on hardware (B200, CUDA 12.8):

• On Hopper/Blackwell the default cudaLimitPersistingL2CacheSize is already nonzero (23.7 MiB out of a 126 MiB L2 on the B200 I tested, in a fresh context). So a submission doesn't even need the cudaDeviceSetLimit call — just setting a persisting access-policy window is enough to start parking data in the set-aside region.
• The benchmark stream's access-policy window is never cleared by the harness, so a window a submission sets stays armed for the whole run and re-pins on every launch.

This matches documented CUDA L2-management semantics (the "Device Memory L2 Access Management" section of the CUDA C++ Programming Guide): the set-aside region is reserved for persisting accesses and is not available to normal/streaming accesses, and releasing it requires resetting cudaLimitPersistingL2CacheSize and/or calling cudaCtxResetPersistingL2Cache(). I verified the consequence directly on the B200 below: a plain spam-write cannot evict a set-aside persisting buffer, but releasing the carveout makes it fully cold again.

The fix

Before the spam-write, clear_cache() now also:

1. Releases the persisting carveout (cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, 0)) so the spam-write uses the whole L2 and can evict everything — including lines the previous kernel parked as persisting. Changing this device limit drains in-flight work, which also guarantees the previous (untrusted) kernel has finished before we demote+evict its lines. This is gated on the carveout actually being nonzero, so honest kernels pay nothing: they never reserve a carveout, so after the first clear the limit stays 0 and the branch is skipped. cudaDeviceGetLimit is cheap and non-blocking (~0.8 µs, measured).
2. Keeps the existing cudaCtxResetPersistingL2Cache() to demote any still-persisting lines.
3. Clears the stream's access-policy window so the next launch doesn't immediately re-mark its inputs persisting.

The change is localized to clear_cache(), which is the single choke point used for both warmup and the timed loop.

Honesty about scope / severity

I want to be upfront: I could not turn this into a working end-to-end score inflation on the B200 I had access to. End-to-end through do_bench_isolated, a warm-resident buffer (none-clear) measured the same as a cold one even for a 124 µs, L2-bandwidth-bound kernel — the B200's HBM3e is fast enough relative to its very large L2 that the warm/cold gap is small, and the stock reset + spam-write does in practice evict on this driver/arch.

So this PR is defense-in-depth hardening of a defense that is fragile by construction, not a patch for a live B200 exploit:

• The current defense's correctness rests entirely on a single host-immediate, non-stream-ordered cudaCtxResetPersistingL2Cache() landing effectively inside a fully pipelined loop that never synchronizes. I verified (a) that call does not synchronize the device (returns in ~5 µs while a ~1 s kernel runs) and (b) with the device idle, a plain spam-write evicts only ~50–65% of a set-aside persisting buffer — the reset is doing the rest, unordered.
• The magnitude of the advantage this guards against scales with the GPU's L2:HBM bandwidth ratio, which is much higher on the A100/H100 (sm_80;90, the wheels this repo ships) and on the L4 used in this repo's CI than on a B200. The fix removes the reliance on reset ordering entirely on those parts.

Microbenchmark evidence (B200, device-idle, isolates the eviction physics)

A standalone test that pins a reused buffer persisting, primes it L2-resident, applies a clear, then times one bandwidth-bound read (warm/cold gap ~1.7x at 64 MiB on this B200):

buf=64MiB | warm(pinned,resident)=12.80us  cold(normal,evicted)=22.18us
  persist + spam-write only        : still ~warm   (carveout shields the lines)
  persist + reset + spam-write     : evicts only with the *unordered* host reset doing the work
  persist + setLimit(0)+reset+win  : 22.30us  -> 101% of warm->cold  (fix evicts a previously-pinned buffer)

So the spam-write alone cannot evict a set-aside persisting buffer; releasing the carveout makes it cold-cache again, independent of reset ordering.

Test

test/l2_persistence.py benchmarks the same memory-bound GEMV honestly and with the persistence trick, and asserts the trick yields no warm-cache speedup (hacked median ≥ 0.88 × honest median). It's a back-to-back relative comparison on whatever GPU runs it, so it's most discriminating on the CI L4. Passes with this PR; correctness controls (submission_correct / submission_wrong in exploits/) are unaffected (verified: correct → PASS, wrong → caught).

Notes

• Built and tested locally with CUDA 12.8 on a B200 (-DCMAKE_CUDA_ARCHITECTURES=90;100). The new code in clear_cache() uses only stable CUDA runtime APIs (cudaDeviceGetLimit / cudaDeviceSetLimit / cudaCtxResetPersistingL2Cache / cudaStreamSetAttribute) plus <cstddef>/<cstring>, so it should build unchanged under the CI's CUDA 13.0 / CUDAARCHS=80;90;100f;120f.
• Heads up, unrelated to this PR: a fresh CUDA 12.8 build of master fails to compile csrc/check.cu, because cuda::ceil_div there is now called with mismatched argument types ((size_t, int)) which no longer deduce. I worked around it locally to build/test; the 12.8 CI row is currently commented out so this PR doesn't address it, but happy to send a separate one-liner.

[brycelelbach]

Human Bryce here: I've observed this class of reward hacking in two separate autoresearch experiments of my own (one on CCCL's histogram algorithm and one on craftax.cu). I asked an agent to check whether this was exploitable with pygpubench, thus this PR.

The hole it describes matches my understanding of the issue it encountered in our runs, and the fix matches the one that I have in my codebase/harnesses.