Provide an option to reset persisting-L2 state between benchmark cells

[robobryce]

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.

[oleksandr-pavlyk]

@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);, where env carries NVBench stream, memory pool, and, possibly, tuning policy.

I did not find cuda::annotated_ptr, cudaStreamSetAttribute, cudaAccessProperty*, or cudaDeviceSetLimit in the CUB histogram benchmark/implementation path from that PR ref. Also, TUNE_MEM_PREFERENCE there is not L2 persistence; it maps to CUB’s BlockHistogramMemoryPreference (GMEM, SMEM, BLEND), meaning privatized histogram storage strategy.

So the cleanup in that commit looks defensive against ambient CUDA context/stream state, not against a direct mutation performed by
DeviceHistogram::HistogramEven. The commit message seems to confirm this deduction.

If right, the defensive reset should perhaps be done before state.exec call 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]

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.