Provide an option to reset persisting-L2 state between benchmark cells #458
Description
Problem
Benchmarks that enable a CUDA persisting-L2 access policy can leave process-wide state behind for later NVBench benchmark cells. In particular, cudaLimitPersistingL2CacheSize and persisting cache lines can survive after the benchmark that configured them. A later benchmark in the same process can then measure a different cache configuration than it requested.
CCCL's DeviceHistogram benchmark work initially handled this by repeating the following before every measured launch:
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, 0);
cudaCtxResetPersistingL2Cache();That is easy for individual benchmarks to forget and requires the custom timer API to keep the reset outside the timed region.
Proposed behavior
NVBench should provide a benchmark option that restores the default persisting-L2 reservation and demotes persisting lines before a measured launch. This could be enabled by default with an opt-out, or exposed as an explicit state/benchmark option.
The reset should happen outside the measured interval and should report CUDA failures rather than silently continuing.
Motivation
Centralizing this in NVBench prevents one benchmark from contaminating later cells and avoids copying the same setup into every affected benchmark. This came up during review of NVIDIA/cccl#10555.
Activity
oleksandr-pavlyk commentedon Aug 11, 2026
@robobryce I am trying to understand the hazard we are protecting against with this request. The commit where these guards are present is NVIDIA/cccl@2ea754e
The launchable times
cub::DeviceHistogram::HistogramEven(..., env);, whereenvcarries NVBench stream, memory pool, and, possibly, tuning policy.I did not find
cuda::annotated_ptr,cudaStreamSetAttribute,cudaAccessProperty*, orcudaDeviceSetLimitin the CUB histogram benchmark/implementation path from that PR ref. Also,TUNE_MEM_PREFERENCEthere is not L2 persistence; it maps to CUB’sBlockHistogramMemoryPreference (GMEM, SMEM, BLEND), meaning privatized histogram storage strategy.So the cleanup in that commit looks defensive against ambient CUDA
context/streamstate, not against a direct mutation performed byDeviceHistogram::HistogramEven. The commit message seems to confirm this deduction.If right, the defensive reset should perhaps be done before
state.execcall instead of being done inside the timed body before every launch.Please clarify. Understanding the hazard to isolate from contributes to making better design choices.
brycelelbach commentedon Aug 11, 2026
This is to prevent reward hacking when agents are working on nvbench benchmarks. A very common reward hack is to use L2 persistence to cheat in between iterations.