PERKS

1. Overview & Why It Is In This Survey

PERKS (Persistent Kernels) is a GPU execution model for iterative, memory-bound applications: Jacobi-style stencil solvers and the conjugate-gradient (CG) loop. Each time step re-reads the same producer→consumer working set from HBM, computes a new field, and writes it back, only to read it again next step. PERKS' insight is that a stencil kernel leaves the SM's registers and shared memory largely idle, and a persistent-thread grid can hold the inter-iteration working set on-chip across all time steps, eliminating the repeated HBM round-trips. The headline result is up to ~2.12× throughput on 2D stencils on an A100.

1.1 Why now: the bandwidth wall and a decade of idle on-chip storage

GPU peak FLOPs have grown far faster than HBM bandwidth for a decade, so a widening class of kernels is bandwidth-bound: the SMs sit waiting on memory, not arithmetic. Iterative stencils and CG are the textbook case — low arithmetic intensity, large fields, hundreds of steps. At the same time, each Ampere SM ships ~256 KB of registers and ~164 KB of shared memory that a stencil kernel barely touches. The opportunity PERKS seizes is the gap between an under-fed memory system and an over-provisioned, idle on-chip storage system: move the recurring working set up the hierarchy into the storage that is already there and unused. The reason this is newly practical is the maturation of cooperative-grid launch (a grid-wide barrier with all blocks co-resident), which lets the iteration loop live inside one kernel so register/shared-memory contents survive across steps.

1.2 Positioning vs prior locality techniques

1.3 Core contributions

2. The Problem: Re-Streaming the Same Bytes Every Time Step

An iterative stencil computes a new field from the old one, repeated for many time steps. In the conventional “one launch per iteration” model, every step launches a fresh global kernel that reads the entire grid from HBM, computes, and writes it back. The arithmetic per element is tiny (a weighted sum of neighbors), so the kernel is firmly bandwidth-bound: its runtime is governed by HBM traffic, not by the SMs' FLOP capacity.

The waste is structural. The field that step $t$ writes is exactly what step $t{+}1$ reads. If a fixed portion of that field could stay on-chip, the system would skip a full HBM write–read round-trip per step. With $N$ time steps over a domain of $D$ bytes, the conventional model moves on the order of $2 N D$ bytes through HBM; PERKS aims to move a small constant of that.

2.1 The memory hierarchy and where reuse can live

An A100 SM exposes a deep on-chip hierarchy. The crucial design choice in any reuse system is which level the recurring working set is parked in. PERKS bets on the two fastest, software-managed levels (registers, shared memory). The L2 cache — the only level CUDA lets you bias toward persistence via an access-policy window — sits below those, shared across all SMs, and is the level our survey is asking about.

2.2 A worked HBM-traffic estimate

Make the waste concrete. Take an fp32 field of $W \times H = 8192 \times 8192$ cells; that is one copy of $D = 4 \cdot 8192^2 \approx 256\,\text{MB}$. Over $N = 1000$ time steps, the baseline reads and writes the whole field each step, so HBM moves about $2 N D \approx 512\,\text{GB}$. If the persistent grid keeps the interior on-chip and only exchanges halos, the HBM traffic collapses to roughly $2 D$ (load once, store the final field) plus a per-step boundary term proportional to the tile perimeter, which for large square tiles is asymptotically negligible. At ~2 TB/s of A100 HBM bandwidth, the baseline's ~512 GB is a hard floor of about a quarter-second of pure memory time; removing it is exactly the headroom PERKS converts into ~2× throughput.

2.3 Why the obvious baselines do not solve it

Three “just do X” strawmen are worth dismissing up front, because PERKS' value is precisely that none of them work.

3. Mechanism: One Persistent Grid That Holds the Working Set

3.1 Traditional launch vs persistent launch

The baseline path is compiled under TRADITIONLAUNCH: a host-side loop that issues one global kernel launch per time step. Between launches there is no guarantee that any register or shared-memory content survives, so the field must be flushed to and reloaded from HBM. The PERKS path is compiled under PERSISTENTLAUNCH: one cooperative launch (cudaLaunchCooperativeKernel) per timed run, not per time step — the time-step loop lives inside the device kernel and runs with a cooperative grid::sync() between steps, so the working set stays on-chip across steps. The outer RUNS loop simply repeats the timed experiment.

Both branches are visible side by side in the source (file paths are real, line numbers are from the cloned repo):

// stencil/2dstencil/jacobi-2d.cu  (~L754–L770)
#ifdef PERSISTENTLAUNCH          // PERKS: one cooperative launch for the WHOLE loop
  for(int i=0; i<RUNS; i++) {
    cudaLaunchCooperativeKernel((void*)execute_kernel,
            executeGridDim, executeBlockDim,
            ExecuteKernelArgs, executeSM, 0);   // iteration loop lives INSIDE the kernel
  }
#endif
#ifdef TRADITIONLAUNCH           // baseline: one global launch PER iteration
  for(int i=0; i<RUNS; i++) {
    execute_kernel<<<executeGridDim, executeBlockDim, executeSM>>>
            (input, width_y, width_x, __var_2__);    // field round-trips through HBM each step
  }
#endif

3.2 Register fold vs shared-memory fold

PERKS exposes two knobs at run time. --blkpsm sets the number of thread blocks resident per SM (and thus how much of the working set each block can cache), and --usesm switches on the shared-memory cache in addition to the register cache. The gen binary caches into registers; gen --usesm adds the SRAM tier; genwr caches into the maximum number of registers first and spills to shared memory. The choice of how many register “folds” vs shared-memory folds to use is precomputed per stencil type and precision in perksconfig.cuh, so a default build needs no codegen.

3.3 Boundary exchange without leaving the chip

A stencil's update at a tile edge needs neighbor values that belong to an adjacent block. In the baseline model that data is communicated implicitly: every block writes its whole tile to HBM, and the next step's launch reads neighbors back. PERKS cannot afford that round-trip — the entire point is to keep the field on-chip. Instead it relies on the persistent grid: all blocks are co-resident, so after each step a grid-wide synchronization makes the just-computed halo cells globally visible (through a thin slab of global memory used only for the boundary, plus the L2_cache3/L2_cache4 scratch buffers), and the next step proceeds. The bulk interior of every tile avoids the repeated HBM round-trip; only boundary / global scratch and the final output still move through HBM. This is why a cooperative launch is mandatory and why occupancy, not arithmetic, becomes the binding constraint.

3.4 The output line we will measure

Binaries built with -D__PRINT__ emit a single tab-separated line per run. The headline number is throughput in GCells/s. From the source, the timing/throughput print is:

// stencil/2dstencil/jacobi-2d.cu  (~L812)  —  elapsedTime(ms) \t throughput(GCells/s)
printf("%f\t%f\t", elapsedTime,
       (REAL)iteration*(width_y-2*HALO)*(width_x-2*HALO) / (elapsedTime/RUNS)/1000/1000);

The speedup we report is simply gen (or genwr) GCells/s divided by baseline GCells/s, for the same stencil type, domain size, and iteration count.

4. The Crux: What PERKS Deliberately Does NOT Use

This is the section that matters for our survey. PERKS' speedup is real, but the question we care about is which level of the hierarchy delivers it. Three independent pieces of evidence — two from the paper, one from the committed source — converge on the same answer: not the L2 cache, and not kernel-launch elision.

4.0 Background: what the Ampere L2 set-aside actually does

Ampere (sm_80) introduced a software-controllable L2 persistence feature. A program can carve out a portion of the ~40 MB L2 as a “set-aside” region and attach an access-policy window to a stream: accesses to a chosen address range get a hitRatio (the fraction of accesses biased toward persistence) and a hitProp of cudaAccessPropertyPersisting, so the hardware preferentially keeps those lines resident instead of evicting them under normal LRU pressure. This is the closest a CUDA programmer can come to “pinning” an array in L2. It is exactly the mechanism a naive reuse story would reach for — and exactly the mechanism PERKS' authors wired up, set to maximum (hitRatio = 1), measured, and discarded.

4.1 Evidence #1 (source): the L2 set-aside path is committed, then commented out

In jacobi-2d.cu, the entire CUDA L2-residency control path — the access-policy window, hitRatio = 1, cudaAccessPropertyPersisting, and cudaCtxResetPersistingL2Cache() — is present in the repository but fully commented out. It is gated behind #if defined(GEN) && defined(L2PER), and no CMake target defines L2PER, so it is dead code in every shipped build. Directly above it is the author's verdict. This is a verifiable, static fact in the cloned repo (confirmed at lines ~715–729):

// stencil/2dstencil/jacobi-2d.cu  L715–L729  (verbatim, confirmed in cloned repo)
//@NEO: l2 cache setting not show performance difference.
#if defined(GEN) && defined(L2PER)
    // REAL l2perused;
    // size_t inner_window_size = 30*1024*1024;
    // cudaStreamAttrValue stream_attribute;
    // stream_attribute.accessPolicyWindow.base_ptr  = reinterpret_cast<void*>(L2_cache3);
    // stream_attribute.accessPolicyWindow.num_bytes = min(inner_window_size, L2_utage_2*2);
    // stream_attribute.accessPolicyWindow.hitRatio  = 1;        // pin in L2
    // stream_attribute.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;
    // stream_attribute.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;
    // cudaStreamSetAttribute(0, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);
    // cudaCtxResetPersistingL2Cache();
    // cudaStreamSynchronize(0);
#endif

Read this carefully: the author wrote the full Ampere L2-persistence API call, set hitRatio = 1 (the maximum “keep this in L2” hint), measured it, found no benefit, commented it out, left a note, and shipped the SRAM/register version instead. This is not an oversight. It is a recorded negative result baked into the code.

4.2 Evidence #2 (paper Sec 4.4.2): forcing L2 residency on A100 gives ~8% slowdown

The paper's own ablation (Section 4.4.2) reports that enabling the A100 L2 residency control produced approximately an 8% slowdown rather than a speedup. The intuition: L2 is shared by all SMs; reserving a persistence window for the stencil field reduces the L2 capacity available to everything else (boundary exchanges, write-backs, instruction/constant traffic), and the field is anyway already cached more cheaply upstream in registers and shared memory. Biasing L2 toward an array that is already resident on-chip is pure overhead.

4.3 Evidence #3 (paper Sec 2.1.2): launch overhead is negligible

One might object that the persistent grid wins by eliding N−1 kernel launches. Section 2.1.2 rules this out: kernel-launch overhead is negligible relative to the kernels' runtime at these problem sizes. The single launch matters only because it is what enables on-chip residency across steps — not because launches themselves were expensive. So the speedup is not a launch-overhead artifact, and it is not an L2 artifact; by elimination it is the register/shared-memory residency.

5. Repository Map & Build Inventory

The repository is github.com/neozhang307/PERKS. We cloned it into the survey at papers/perks/repo and on PACE at $PROJECT/repro/PERKS. The relevant tree for the 2D-stencil reproduction:

5.1 Confirmed static facts from the source

5.2 The conjugate-gradient variant (optional, off the headline path)

The repository also ships a conjugate-gradient solver under conjugateGradient/ that applies the same persistent-grid idea to the CG iteration (the SpMV + vector-update loop reuses the sparse operands across steps). It builds with the same config.sh / build.sh pattern but requires --mtx matrix inputs (e.g. SuiteSparse files), which adds data-staging overhead and varies the result by matrix. Because our claim is about the activation/array-reuse term, the headline reproduction stays on the 2D stencil, where the only reusable bytes are the field arrays and there is no sparse-format or input-data confound. CG is a cheap add if time permits but is not on the critical path for the L2 conclusion.

5.3 Build (completed on PACE)

The 2D-stencil build completed on an A100 node; the per-type binaries land in build/init/<type>/ as <type>_baseline.exe, <type>_gen.exe, <type>_genwr.exe, plus naive/reference variants. Build needs only the CUDA toolkit, CMake ≥ 3.10, and a C++14 host compiler — no exotic dependencies.

# PACE build (completed): per-component cmake
cd $PROJECT/repro/PERKS/stencil/2dstencil
bash config.sh   # rm -f build/init/CMakeCache.txt ; cmake -S . -B ./build/init/
bash build.sh    # cmake --build ./build/init/ -j 16
# -> build/init/2d9pt/2d9pt_baseline.exe , _gen.exe , _genwr.exe  (fat sm_70;sm_80)

5.4 Claims vs implementation

6. Our Reproduction (A100-80GB) — Measured on PACE (job 9934117)

6.1 Hypothesis & outcome

Our hypothesis going in was: on a large fp32 2D stencil at W=H=8192, iter=1000, bdim=256, the PERKS gen / genwr binary will reach roughly 1.8–2.2× the baseline GCells/s, bracketing the paper's 2.12× suite figure for a single representative stencil. We further expected that the speedup is driven by register + shared-memory residency rather than higher L2 reuse, so the L2 hit rate should stay roughly flat — and that when L2 residency is forced on (our -DL2PER rebuild), throughput should drop by roughly the paper's ~8%. The reasoning is the elimination argument of Section 4: launch overhead is negligible (Sec 2.1.2), the field is staged on-chip upstream of L2, and forced L2 persistence is net-negative (Sec 4.4.2, plus the commented-out code).

Outcome (measured, job 9934117): the throughput prediction landed — full PERKS (genwr) measured 1.86× (2d5pt), 1.90× (2d9pt), 2.13× (2d13pt) vs baseline, confirming the paper's ~2.12× geomean. The L2-attribution prediction landed too: L2 sector hit-rate stayed essentially flat (baseline ≈72% → PERKS ≈79%, only ~7 pp) while DRAM traffic fell sharply, so the byte savings come from on-chip register/shared-memory staging upstream of L2, not from a higher L2 hit rate. The decisive -DL2PER residency rebuild is the one piece we did NOT obtain this run (cmake reconfigure failed; see Sec 7.3) — that single leg rests on the paper's Sec 4.4.2 (~8% slowdown) plus the committed-but-commented code, not on our own timing.

6.2 Method (exactly what we build & run)

Single A100-80GB, exclusive job (gpu-a100, --gres=gpu:a100:1, --constraint=A100-80GB, paid account). A100 is the only faithful target: the code is hard-locked to sm_70/sm_80 in CMake and the runtime ptx dispatch only branches on 700/800, so L40S (sm_89) and H100/H200 (sm_90) would mis-dispatch or fail to build — there is no meaningful multi-arch fan-out. We run the binaries directly with CUDA_VISIBLE_DEVICES=0 and explicitly DROP the authors' iterategen.sh / temp.sh, which assume a 2-GPU node and gate every run on nvidia-smi -i 1 temperature.

# 1) Timed reproduction (per stencil type in {2d5pt, 2d9pt, 2d13pt})
W=8192; H=8192; ITER=1000; BDIM=256
# correctness first (mirrors the authors' --check at iter=3)
2d9pt_gen.exe      $W $H --warmup --iter=3    --bdim=$BDIM --blkpsm=1 --fp32 --check
# baseline (no inter-iteration caching)
2d9pt_baseline.exe $W $H --warmup --iter=$ITER --bdim=$BDIM --blkpsm=0 --fp32
# PERKS: register cache / register+SRAM cache / write-spill
2d9pt_gen.exe      $W $H --warmup --iter=$ITER --bdim=$BDIM --blkpsm=1 --fp32
2d9pt_gen.exe      $W $H --warmup --iter=$ITER --bdim=$BDIM --blkpsm=1 --usesm --fp32
2d9pt_genwr.exe    $W $H --warmup --iter=$ITER --bdim=$BDIM --blkpsm=1 --fp32
# speedup = gen|genwr GCells/s  /  baseline GCells/s

6.3 Results — measured on PACE A100 (job 9934117)

n/r = not separately recorded in our measurement log. The decisive intermediate stage (gen / gen --usesm) is recorded only for the representative type 2d9pt, where gen (register cache) is 1.40× and gen --usesm (register+SRAM) is 1.34× over baseline; the per-stencil baseline and full-PERKS (genwr) throughputs and the speedup column are the recorded measured values.

7. L2 Attribution — Isolating the Reuse Level with Nsight Compute

The reproduction shows that PERKS is faster; the attribution shows why, and crucially decomposes the speedup into the levels of the hierarchy. We profiled baseline vs gen (and gen --usesm) with Nsight Compute on a representative type (2d9pt) so the ncu kernel replay is cheap. The intended decisive separate experiment is the -DL2PER rebuild; this run obtained the L2-hit-rate and DRAM measurements but not the -DL2PER timing (its cmake reconfigure failed — see Sec 7.3).

7.1 The metric set (sm_80-valid)

# Nsight Compute, one representative type (2d9pt), iter=30, 4 launches
ncu --target-processes all --launch-count 4 -f -o ncu_gen \
    --metrics dram__bytes_read.sum,dram__bytes_write.sum,\
lts__t_sector_hit_rate.pct,lts__t_sectors_lookup_hit.sum,lts__t_sectors_lookup_miss.sum,\
l1tex__data_pipe_lsu_wavefronts_mem_shared.sum,\
smsp__inst_executed_op_shared_ld.sum,smsp__inst_executed_op_shared_st.sum,\
launch__grid_size,sm__cycles_elapsed.avg \
    2d9pt_gen.exe $W $H --iter=30 --bdim=256 --blkpsm=1 --fp32

7.2 The attribution logic

7.3 Attribution results — measured on PACE A100 (job 9934117)

8. Scope, Threats to Validity & Transfer to LLM Inference

8.1 Reproduction risks we are tracking

9. Takeaway

PERKS is a high-quality systems paper about doing inter-iteration reuse right on the GPU. For our survey its value is almost entirely in what it refuses to do: a team of GPU-memory specialists, building the most reuse-aggressive stencil engine they could, deliberately staged the recurring working set in registers and shared memory, and explicitly rejected the L2 cache — both in prose (Sec 4.4.2: ~8% slowdown) and in code (the commented-out access-policy-window block with the “no performance difference” note). Our reproduction made that decision quantitative on a single A100 (job 9934117): we measured full PERKS at 1.86–2.13× over baseline (confirming the paper's ~2.12× geomean), and we measured DRAM traffic falling sharply while the L2 sector hit-rate stayed essentially flat (≈72% → ≈79%, ~7 pp) — the signature that the byte savings come from on-chip register/shared-memory staging, not from L2. The one leg we did NOT obtain ourselves is the forced-L2 (-DL2PER) timing: that rebuild patched the source but its cmake reconfigure failed this run, so the ~8% slowdown rests on the paper's Sec 4.4.2 plus the committed-but-commented code, and is flagged for a follow-up run. On the two legs we did measure, PERKS becomes a clean external corroboration — qualitative, on a non-LLM iterative-stencil workload, not a direct transfer of the numeric bound — of our “Worth the Bytes?” thesis: L2-mediated transient reuse is not the win, the gain is not driven by higher L2 reuse, and an explicit L2-residency policy did not pay — consistent with our <0.65% bound.