From f9fb5a007ffaadd2926e456b884291ce6a449138 Mon Sep 17 00:00:00 2001 From: Bryce Adelstein Lelbach Date: Sun, 7 Jun 2026 16:19:40 +0000 Subject: [PATCH] Harden L2 cache clearing against persistence reward-hacking clear_cache() spam-writes 2*L2 of dummy data to evict the previous iteration's footprint, and calls cudaCtxResetPersistingL2Cache() to undo "sneaky cache persistency". That reset alone is not sufficient when a submission also reserves a persisting L2 carveout: a kernel can mark a reused buffer (weights, lookup tables, ...) persisting via a stream access-policy window (cudaStreamSetAttribute / cudaStreamAttributeAccessPolicyWindow, hitProp=cudaAccessPropertyPersisting) backed by cudaLimitPersistingL2CacheSize. Normal/streaming writes can only use the *non*-set-aside portion of L2, so the spam-write cannot evict lines parked in the carveout -- the buffer stays warm across timed iterations and the kernel is scored as if its inputs were always L2-resident. On Hopper/Blackwell the default persisting carveout is already nonzero, so a submission only needs to set an access-policy window (no cudaDeviceSetLimit call) to start parking data in the set-aside region. Before the spam-write, also: - release the persisting carveout (cudaDeviceSetLimit(..., 0)) so the write uses the whole L2 and can evict everything, including previously-parked persisting lines. This is gated on the carveout actually being nonzero, so honest kernels (which never reserve one) pay nothing after the first clear; cudaDeviceGetLimit is cheap and non-blocking. - clear the stream's access-policy window so the next launch does not immediately re-mark its inputs persisting. Adds test/l2_persistence.py, which benchmarks a memory-bound GEMV honestly and with the persistence trick and asserts the trick confers no speedup. The guarded advantage scales with the GPU's L2:HBM bandwidth ratio (large on A100/H100/L4), so it is a back-to-back relative comparison on the same device. Co-Authored-By: Claude Opus 4.8 (1M context) --- csrc/clear_l2.cu | 39 ++++++++++++- test/l2_persistence.py | 126 +++++++++++++++++++++++++++++++++++++++++ 2 files changed, 164 insertions(+), 1 deletion(-) create mode 100644 test/l2_persistence.py diff --git a/csrc/clear_l2.cu b/csrc/clear_l2.cu index ece134f..f7bd907 100644 --- a/csrc/clear_l2.cu +++ b/csrc/clear_l2.cu @@ -2,6 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // +#include +#include #include #include @@ -24,9 +26,44 @@ __global__ void discard_cache_kernel(uint4* dummy_memory, const int size) { :"memory"); } +// Neutralize any L2 *persistence* setup before the spam-write, so the write can actually +// evict everything. A submission can otherwise keep a buffer resident in L2 across timed +// iterations by marking it persisting via an access-policy window +// (cudaStreamSetAttribute, hitProp=cudaAccessPropertyPersisting) backed by a set-aside L2 +// carveout (cudaLimitPersistingL2CacheSize). Streaming/normal writes can only use the +// *non*-set-aside portion of L2, so the write below cannot evict lines parked in the +// carveout -- the warm buffer then survives the "cache clear" and inflates the score. +// +// Note: on recent GPUs (e.g. Hopper/Blackwell) the default persisting carveout is already +// nonzero, so a submission only needs to set an access-policy window -- no explicit +// cudaDeviceSetLimit call -- to start parking data in the set-aside region. +static void disarm_l2_persistence(cudaStream_t stream) { + // Release the set-aside carveout so the spam-write can use the whole L2 and 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. Only pay that cost when a + // carveout is actually reserved: honest kernels never set it, so after the first clear + // the limit stays 0 and this branch is skipped (cudaDeviceGetLimit is cheap/non-blocking). + std::size_t persist_limit = 0; + if (cudaDeviceGetLimit(&persist_limit, cudaLimitPersistingL2CacheSize) == cudaSuccess && + persist_limit != 0) { + CUDA_CHECK(cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, 0)); + } + + // Demote any lines still flagged persisting to normal status so the write can evict them. + CUDA_CHECK(cudaCtxResetPersistingL2Cache()); + + // Drop any access-policy window the submission left on the benchmark stream, so the next + // kernel launch does not immediately re-mark its inputs persisting and re-pin them. + cudaStreamAttrValue window_attr; + std::memset(&window_attr, 0, sizeof(window_attr)); + window_attr.accessPolicyWindow.num_bytes = 0; + CUDA_CHECK(cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &window_attr)); +} + void clear_cache(void* dummy_memory, int size, bool discard, cudaStream_t stream) { // make sure there's no sneaky cache persistency setting that defeats our spam-writing - CUDA_CHECK(cudaCtxResetPersistingL2Cache()); + disarm_l2_persistence(stream); int nelem = size / sizeof(uint4); // write a large amount of memory to ensure all cache lines are cleared diff --git a/test/l2_persistence.py b/test/l2_persistence.py new file mode 100644 index 0000000..9456168 --- /dev/null +++ b/test/l2_persistence.py @@ -0,0 +1,126 @@ +"""L2-persistence reward-hacking regression test. + +A submission can try to inflate its score by keeping a large, *reused* input (weights, +lookup tables, ...) resident in L2 across the benchmark's timed iterations, so the timed +kernel reads it from warm L2 instead of from HBM. It does this by marking the buffer +persisting with a stream access-policy window (``cudaStreamSetAttribute`` / +``cudaStreamAttributeAccessPolicyWindow``, ``hitProp=cudaAccessPropertyPersisting``) backed +by a set-aside L2 carveout (``cudaLimitPersistingL2CacheSize``). Plain streaming writes can +only use the *non*-set-aside portion of L2, so the harness's L2-clearing spam-write cannot +evict the parked lines -- unless the harness first releases the carveout and demotes the +persisting lines (see ``csrc/clear_l2.cu``). On Hopper/Blackwell the default persisting +carveout is already nonzero, so a submission only needs the access-policy window. + +The test runs the SAME memory-bound GEMV through the isolated benchmark twice: + * ``kernel`` -- honest (each iteration reads the weights from a cleared L2). + * ``kernel_l2_persist`` -- marks the reused weights persisting so they would survive the + clear and stay warm across iterations. +With the clear_cache hardening, the persisting variant gets no warm-cache advantage: its +time is the same as the honest one (within noise). The size of the advantage this guards +against scales with the GPU's L2:HBM bandwidth ratio -- large on A100/H100/L4, small on +parts with very fast HBM -- so the assertion is a relative back-to-back comparison on the +same device, and is most discriminating on the L4 used in CI. +""" +import ctypes +import torch +import pygpubench + +_dev = torch.cuda.current_device() +_L2 = torch.cuda.get_device_properties(_dev).L2_cache_size +_cudart = ctypes.CDLL("libcudart.so") + +# cudaDevAttrMaxPersistingL2CacheSize == 108; fall back to half of L2 if unavailable. +_max_persist = _L2 // 2 +_v = ctypes.c_int(0) +if _cudart.cudaDeviceGetAttribute(ctypes.byref(_v), 108, _dev) == 0 and _v.value > 0: + _max_persist = _v.value + +# Reused weight matrix sized to fit in the persisting carveout (cap 64 MiB). +_W_BYTES = min(int(_max_persist * 0.75), 64 * 1024 * 1024) +_M = 2048 +_K = max(1, _W_BYTES // (4 * _M)) +_W = torch.ones(_K, _M, device="cuda", dtype=torch.float32).contiguous() + + +class _AccessPolicyWindow(ctypes.Structure): + _fields_ = [("base_ptr", ctypes.c_void_p), ("num_bytes", ctypes.c_size_t), + ("hitRatio", ctypes.c_float), ("hitProp", ctypes.c_int), ("missProp", ctypes.c_int)] + + +class _StreamAttrValue(ctypes.Union): + _fields_ = [("accessPolicyWindow", _AccessPolicyWindow), ("_pad", ctypes.c_byte * 64)] + + +_CUDA_ACCESS_PROPERTY_STREAMING = 1 +_CUDA_ACCESS_PROPERTY_PERSISTING = 2 +_CUDA_STREAM_ATTR_ACCESS_POLICY_WINDOW = 1 +_CUDA_LIMIT_PERSISTING_L2 = 0x05 +_installed = False + + +def _install_persistence(): + """Pin _W in L2 by reserving a carveout and setting a persisting access-policy window.""" + nbytes = _W.numel() * _W.element_size() + want = min(_max_persist, nbytes) + _cudart.cudaDeviceSetLimit(ctypes.c_int(_CUDA_LIMIT_PERSISTING_L2), ctypes.c_size_t(want)) + stream = torch.cuda.current_stream().cuda_stream + val = _StreamAttrValue() + val.accessPolicyWindow.base_ptr = ctypes.c_void_p(_W.data_ptr()) + val.accessPolicyWindow.num_bytes = want + val.accessPolicyWindow.hitRatio = 1.0 + val.accessPolicyWindow.hitProp = _CUDA_ACCESS_PROPERTY_PERSISTING + val.accessPolicyWindow.missProp = _CUDA_ACCESS_PROPERTY_STREAMING + _cudart.cudaStreamSetAttribute(ctypes.c_void_p(stream), + ctypes.c_int(_CUDA_STREAM_ATTR_ACCESS_POLICY_WINDOW), + ctypes.byref(val)) + + +def kernel(output, x): + torch.mv(_W, x, out=output) + + +def kernel_l2_persist(output, x): + global _installed + if not _installed: + _install_persistence() + _installed = True + torch.mv(_W, x, out=output) + + +def generate_test_case(*, seed): + gen = torch.Generator(device="cuda") + gen.manual_seed(seed) + x = torch.rand(_M, device="cuda", dtype=torch.float32, generator=gen).contiguous() + y = torch.empty(_K, device="cuda", dtype=torch.float32).contiguous() + expected = torch.mv(_W, x).contiguous() + return (y, x), (expected, 1e-2, 1e-2) + + +def _bench(qualname, repeats=120, seed=1234): + # landlock/mseal disabled for portability (e.g. Modal CI); see test/grayscale.py. + res = pygpubench.do_bench_isolated(qualname, generate_test_case, {}, repeats, seed, + discard=True, landlock=False, mseal=False, allow_root=True) + assert res.success, f"{qualname}: benchmark did not succeed (errors={res.errors}, full={res.full})" + return pygpubench.basic_stats(res.time_us) + + +if __name__ == "__main__": + w_mib = _W.numel() * _W.element_size() / (1024 * 1024) + print(f"weights = {w_mib:.0f} MiB (K={_K} x M={_M}), " + f"max_persist = {_max_persist/1048576:.0f} MiB, L2 = {_L2/1048576:.0f} MiB") + + honest = _bench("l2_persistence.kernel") + hacked = _bench("l2_persistence.kernel_l2_persist") + print(f"honest : {honest}") + print(f"l2-persistence : {hacked}") + + # With clear_cache hardened, the persistence trick must confer no warm-cache speedup: + # the attack's median must not beat the honest median beyond run-to-run noise. + ratio = hacked.median / honest.median + print(f"hacked/honest median ratio = {ratio:.3f} (require >= 0.88)") + assert ratio >= 0.88, ( + f"L2-persistence trick gained a warm-cache advantage: hacked median " + f"{hacked.median:.2f}us vs honest {honest.median:.2f}us (ratio {ratio:.3f}); " + f"clear_cache failed to evict the persisting buffer between iterations." + ) + print("OK: L2 persistence confers no benchmark advantage.")