Read-Only Cache: Keep a Lookup Page Hot in L1

Experimental — correctness is graded; the performance target awaits GPU calibration. The profile table has 128 pages: 4 MiB fits in the GPU's shared L2 cache, but not in one SM's L1/TEX cache. The starter assigns round-robin logical queries directly to threads, so every 256-thread block cycles through all 128 pages. Remap workers so each block completes the 256 queries for one 32 KiB page, then write each result back to its original logical output slot.

Languages: CUDA C++, Rust, HIP C++

Lesson

The GPU's shared L2 cache is much larger than the per-SM L1/TEX cache, but an L1 hit is closer to the executing threads and avoids an L2 request. An explicit __ldg load makes the table read-only for this experiment; it does not guarantee that the requested data fits in L1. Cache reuse depends on the active working set. Grouping work by table page lets a block repeatedly reuse a small page instead of cycling through the entire L2-resident table.

Working-set size and reuse determine whether caching helps. Compare repeated paired times on the selected GPU; the speed target is experimental until calibrated.

What you will learn

Distinguish the per-SM L1/TEX cache from the GPU-wide shared L2 cache.
Calculate the starter's 4 MiB per-block table footprint and the goal's 32 KiB page footprint.
Transpose the worker-to-query mapping without changing logical output order.
Interpret higher L1 reuse, lower L2 read traffic, and GPU time together.

Starter code

    int q = blockIdx.x * blockDim.x + threadIdx.x;
    if (q >= N) return;

    int pageCount = N / kQueriesPerPage;
    int page = q % pageCount;
    int localQuery = q / pageCount;
    int pageBase = page * kLookupPageEntries;

    float sum = 0.0f;
    #pragma unroll 1
    for (int lookup = 0; lookup < kLookupsPerQuery; ++lookup) {
        int offset = (localQuery * 131 + lookup * 977)
                & (kLookupPageEntries - 1);
        sum += __ldg(&table[pageBase + offset]);
    }
    output[q] = sum;