From b9adefee0e75dddeedc53bb562c4ae1a11697aa2 Mon Sep 17 00:00:00 2001 From: AndreSlavescu Date: Thu, 30 Jul 2026 12:59:57 -0700 Subject: [PATCH] Correct the memory hierarchy: the earlier table measured the thread count, not the part The hierarchy constants I added a few commits ago were wrong, and an occupancy probe written to explain something else is what showed it. That table put the fast level's capacity at 2 MB with a fourfold cliff past it. There is no cliff at 2 MB. Every size in it was measured at 64 threadgroups, which saturates DRAM at 128 MB and starves an 8 MB working set by a factor of twelve: working set 64 groups 512 groups 8 MB 196 2403 16 MB 160 1910 32 MB 131 447 So the "levels" between 4 and 32 MB were the thread count. Measuring each size at several threadgroup counts and keeping the best moves the capacity from 2 MB to at least 16 MB and turns the far side from a cliff into a ramp: 2453, 1795, 448, 175, 156 GB/s at 16, 32, 64, 128 and 256 MB. The thread count alone moves 16 MB by 19.7x, which is why one count cannot serve the sweep. The DRAM constant survives, and that was worth checking rather than assuming, because the corrected sweep reads 156 GB/s at 256 MB against a recorded 120.6 and a constant that low would have reopened every decode conclusion built on it. At a one-gigabyte working set bandwidth converges to 125 to 131 GB/s whatever the occupancy, so 120.6 is right for true streaming and the 156 was partial residency. STREAMING_READ_GBPS stands. Also recorded, with a caveat that matters more than the numbers: how many threadgroups a resident working set needs before it reads at resident speed, 256 at 4 MB rising to 1024 at 16 MB. Residency is necessary and not sufficient. It is not yet shown to be actionable for real kernels -- the affine decode matmul launches 35 to 280 groups depending on block_n, an eightfold range straddling those thresholds, and its configs measure within 3% of each other, so whatever binds that kernel it is not this. One test changed rather than being made to pass. It asserted the drop past the resident capacity was at least threefold, which encoded the cliff the corrected measurement removed, so it was testing an artefact of the harness. It now asserts the span and the ordering, which is what survives. A second test pins that residency alone does not deliver resident bandwidth. 687 pass. Lint and vulture clean. Co-Authored-By: Claude Opus 5 (1M context) --- benchmarks/agx_memory_hierarchy.py | 99 +++++++++--------- benchmarks/agx_occupancy.py | 158 +++++++++++++++++++++++++++++ metile/target/__init__.py | 2 + metile/target/agx.py | 65 ++++++------ tests/test_target.py | 27 +++-- 5 files changed, 266 insertions(+), 85 deletions(-) create mode 100644 benchmarks/agx_occupancy.py diff --git a/benchmarks/agx_memory_hierarchy.py b/benchmarks/agx_memory_hierarchy.py index ade412b..ce44f71 100644 --- a/benchmarks/agx_memory_hierarchy.py +++ b/benchmarks/agx_memory_hierarchy.py @@ -13,21 +13,20 @@ Three things make the numbers trustworthy rather than suggestive: - saturation the sweep has to reach the known 120.6 GB/s at large sizes, or the thread count is too - low to saturate and every number is a lower bound on something else. Printed as a check - rather than assumed. + occupancy every working set is measured at several threadgroup counts and the best is kept. A fixed + count is not safe: the first version of this used 64 for everything, which saturates DRAM + at 128 MB but starves an 8 MB working set by a factor of twelve, and the resulting table + described the thread count rather than the hierarchy. amortisation each dispatch reads gigabytes and runs for tens of milliseconds, so launch overhead is far below the noise. Timing small kernels through a host round trip is what produced three fabricated results earlier in this project. interleaving sizes are measured in rotating order across rounds, because a sweep that walks from small to large measures thermal drift as much as it measures the hierarchy. -One limit of this probe is worth knowing before reading its output. Inside the resident regime bandwidth -*rises* with working set -- 1545, 2006, 2138 and 2386 GB/s at 256 KB, 512 KB, 1 MB and 2 MB -- which -cannot be a property of a cache. A smaller working set means fewer inner iterations per pass, so the -pass loop's bookkeeping takes a larger share, and the effect shrinks as the set grows. The probe -therefore establishes that a resident set runs at 2386 GB/s or better and says nothing reliable about -how that varies below 2 MB. `metile.target.agx` records it as one number for that reason. +One limit remains after the occupancy sweep. The smallest working sets cannot be given both enough +threads to saturate and enough work per thread to amortise the pass loop, because the two demands +conflict once the set is only a few times the thread count. Numbers below about 1 MB are floors rather +than levels, and `metile.target.agx` records the resident regime as a single figure for that reason. usage: python benchmarks/agx_memory_hierarchy.py @@ -50,11 +49,15 @@ from metile.target import agx THREADGROUP = 256 -# Enough threads to saturate DRAM, and few enough that a 256 KB working set still gives every thread -# something to do: the inner loop strides by the total thread count, so a working set smaller than that -# leaves threads idle and measures parallelism instead of bandwidth. -THREADGROUPS = 64 -VECTOR_BYTES = 16 # float4 +VECTOR_BYTES = 16 + +# Threadgroup counts tried at every working set, with the best taken. A single count cannot serve the +# whole sweep and picking one produced a wrong answer that sat in the target model until the occupancy +# probe contradicted it: at 64 groups an 8 MB working set reads 196 GB/s and at 512 it reads 2403, so the +# "level" recorded at 8 MB was the thread count, not the memory system. Too few threads starves the path; +# too many leaves each thread one inner iteration per pass, where loop bookkeeping competes with the +# loads. Only the maximum over the sweep is a property of the part. +THREADGROUP_COUNTS = (32, 64, 128, 256, 512, 1024) # float4 SIZES_KB = (256, 512, 1024, 2048, 4096, 8192, 16384, 32768, 65536, 131072, 262144) @@ -63,7 +66,6 @@ def _arguments(): parser = argparse.ArgumentParser(description=__doc__) parser.add_argument("--traffic", type=float, default=4.0, help="GB read per dispatch") parser.add_argument("--rounds", type=int, default=5) - parser.add_argument("--threadgroups", type=int, default=THREADGROUPS) return parser.parse_args() @@ -93,26 +95,30 @@ def main(): from metile.runtime.metal_device import MetalDevice device = MetalDevice.get() - threads = arguments.threadgroups * THREADGROUP - grid, block = (threads, 1, 1), (THREADGROUP, 1, 1) pipeline = device.compile_msl(READ_KERNEL, "probe") print(f"device: {device.name}") - print(f"{threads} threads, {arguments.traffic:g} GB read per dispatch") + print(f"{arguments.traffic:g} GB read per dispatch, best of {THREADGROUP_COUNTS} threadgroups") print(f"recorded streaming ceiling: {agx.STREAMING_READ_GBPS} GB/s\n") target_bytes = arguments.traffic * 1e9 - cases = [] largest = max(SIZES_KB) * 1024 data = metile.Buffer(data=np.ones(largest // 4, dtype=np.float32)) - out = metile.Buffer(data=np.zeros(threads * 4, dtype=np.float32)) + out = metile.Buffer(data=np.zeros(max(THREADGROUP_COUNTS) * THREADGROUP * 4, dtype=np.float32)) + cases = [] for size_kb in SIZES_KB: size = size_kb * 1024 vectors = size // VECTOR_BYTES - if vectors < threads: - continue passes = max(1, int(target_bytes // size)) + for count in THREADGROUP_COUNTS: + if vectors < count * THREADGROUP: + continue + cases.append((size, count, vectors, passes)) + + def measure(case): + _, count, vectors, passes = case + threads = count * THREADGROUP buffers = [ data.metal_buffer, out.metal_buffer, @@ -120,48 +126,45 @@ def main(): metile.Buffer(data=np.array([passes], dtype=np.uint32)).metal_buffer, metile.Buffer(data=np.array([threads], dtype=np.uint32)).metal_buffer, ] - cases.append((size, vectors, passes, buffers)) - - def measure(buffers): started = time.perf_counter_ns() - device.dispatch_kernel(pipeline, buffers, grid, block) + device.dispatch_kernel(pipeline, buffers, (threads, 1, 1), (THREADGROUP, 1, 1)) device.sync() - return (time.perf_counter_ns() - started) / 1e9 + seconds = (time.perf_counter_ns() - started) / 1e9 + return vectors * VECTOR_BYTES * passes / seconds / 1e9 - for _, _, _, buffers in cases: - measure(buffers) + for case in cases: + measure(case) - samples = {size: [] for size, _, _, _ in cases} + samples = {case: [] for case in cases} for index in range(arguments.rounds): ordered = cases[index % len(cases) :] + cases[: index % len(cases)] - for size, _, _, buffers in ordered: - samples[size].append(measure(buffers)) + for case in ordered: + samples[case].append(measure(case)) + + rates = {case: statistics.median(values) for case, values in samples.items()} - print(f"{'working set':>13}{'passes':>8}{'ms':>9}{'GB/s':>9}{'vs DRAM':>9}") + print(f"{'working set':>13}{'best groups':>13}{'GB/s':>9}{'vs DRAM':>9}{'worst groups':>14}") results = [] - for size, vectors, passes, _ in cases: - seconds = statistics.median(samples[size]) - gbps = (vectors * VECTOR_BYTES * passes) / seconds / 1e9 - results.append((size, gbps)) + for size_kb in SIZES_KB: + size = size_kb * 1024 + here = [(case, rate) for case, rate in rates.items() if case[0] == size] + if not here: + continue + best_case, best = max(here, key=lambda pair: pair[1]) + worst = min(rate for _, rate in here) + results.append((size, best)) label = f"{size // 1024} KB" if size < 1024 * 1024 else f"{size // (1024 * 1024)} MB" print( - f"{label:>13}{passes:>8}{seconds * 1e3:>9.1f}{gbps:>9.1f}" - f"{gbps / agx.STREAMING_READ_GBPS:>8.2f}x" + f"{label:>13}{best_case[1]:>13}{best:>9.0f}" + f"{best / agx.STREAMING_READ_GBPS:>8.2f}x{best / worst:>13.1f}x" ) dram = min(gbps for _, gbps in results) peak = max(gbps for _, gbps in results) - print(f"\nslowest {dram:.1f} GB/s, fastest {peak:.1f} GB/s, ratio {peak / dram:.2f}x") - saturated = dram >= 0.9 * agx.STREAMING_READ_GBPS + print(f"\nslowest {dram:.0f} GB/s, fastest {peak:.0f} GB/s, ratio {peak / dram:.2f}x") print( - f"saturation check: largest working sets reach {dram:.1f} GB/s against a recorded " - f"{agx.STREAMING_READ_GBPS} GB/s -- {'ok' if saturated else 'TOO LOW, raise --threadgroups'}" + "last column is how much the thread count alone moves that size, which is why it is swept." ) - if not saturated: - print( - "Until that passes, every number here is a lower bound and the knees may be artefacts." - ) - return 1 # Report the knees rather than leaving them to be eyeballed: a level boundary is where bandwidth # drops materially between adjacent sizes. diff --git a/benchmarks/agx_occupancy.py b/benchmarks/agx_occupancy.py new file mode 100644 index 0000000..1c3c1d0 --- /dev/null +++ b/benchmarks/agx_occupancy.py @@ -0,0 +1,158 @@ +"""How much parallelism a kernel needs before memory bandwidth saturates. + +Decode kernels achieve 80 to 128 GB/s while the memory level their working set sits at delivers 128 to +555. That gap is not tiling -- decode reads each weight once, so there is no reuse to capture -- which +leaves the question of what it is. The first suspect is that the kernel is not running enough +threadgroups to saturate the path in the first place. + +That is a property of the part, not of any kernel, so it can be measured directly: hold the working set +fixed and sweep the number of threadgroups. The curve's knee is the parallelism a kernel needs, and any +kernel launching fewer than that is leaving bandwidth behind for a reason a tile-size change could fix. + +The number matters for tile selection specifically. A GEMM tiled by output width launches +`output_features / block_n` threadgroups, so a wider tile means fewer of them: at N=8960 a block_n of 64 +gives 140 threadgroups and 256 gives 35. If saturation needs more than 35, the widest tiles are +self-limiting and the compiler can rule them out before measuring anything. + +Same controls as the other probes here: each dispatch reads gigabytes so launch overhead is negligible, +thread counts are measured in rotating order so drift does not land on one end of the sweep, and elapsed +time has to scale with the pass count or the loop is being collapsed. +""" + +import argparse +import statistics +import sys +import time +from pathlib import Path + +import numpy as np + +_root = str(Path(__file__).resolve().parent.parent) +sys.path.insert(0, _root) + +import metile + +THREADGROUP = 256 +VECTOR_BYTES = 16 +THREADGROUP_COUNTS = (1, 2, 4, 8, 16, 24, 32, 48, 64, 96, 128, 192, 256, 384, 512) + +# Two footprints: one in the resident regime and one streaming, because the parallelism needed to +# saturate a cache and to saturate DRAM are not obviously the same number. +WORKING_SETS = (1024 * 1024, 8 * 1024 * 1024, 64 * 1024 * 1024) + +KERNEL = """#include +using namespace metal; + +kernel void probe(device const float4* data [[buffer(0)]], + device float4* out [[buffer(1)]], + constant uint& vectors [[buffer(2)]], + constant uint& passes [[buffer(3)]], + constant uint& stride [[buffer(4)]], + uint gid [[thread_position_in_grid]]) { + float4 total = float4(0.0f); + for (uint pass = 0; pass < passes; ++pass) { + for (uint index = gid; index < vectors; index += stride) { + total += data[index]; + } + } + out[gid] = total; +} +""" + + +def _arguments(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--traffic", type=float, default=4.0, help="GB read per dispatch") + parser.add_argument("--rounds", type=int, default=3) + return parser.parse_args() + + +def main(): + arguments = _arguments() + from metile.runtime.metal_device import MetalDevice + + device = MetalDevice.get() + pipeline = device.compile_msl(KERNEL, "probe") + largest = max(WORKING_SETS) + data = metile.Buffer(data=np.ones(largest // 4, dtype=np.float32)) + out = metile.Buffer(data=np.zeros(max(THREADGROUP_COUNTS) * THREADGROUP * 4, dtype=np.float32)) + + print(f"device: {device.name}") + print(f"{THREADGROUP} threads per group, {arguments.traffic:g} GB read per dispatch\n") + + def measure(threadgroups, size): + threads = threadgroups * THREADGROUP + vectors = size // VECTOR_BYTES + passes = max(1, int(arguments.traffic * 1e9 // size)) + buffers = [ + data.metal_buffer, + out.metal_buffer, + metile.Buffer(data=np.array([vectors], dtype=np.uint32)).metal_buffer, + metile.Buffer(data=np.array([passes], dtype=np.uint32)).metal_buffer, + metile.Buffer(data=np.array([threads], dtype=np.uint32)).metal_buffer, + ] + started = time.perf_counter_ns() + device.dispatch_kernel(pipeline, buffers, (threads, 1, 1), (THREADGROUP, 1, 1)) + device.sync() + seconds = (time.perf_counter_ns() - started) / 1e9 + return vectors * VECTOR_BYTES * passes / seconds / 1e9 + + cases = [ + (count, size) + for size in WORKING_SETS + for count in THREADGROUP_COUNTS + if size // VECTOR_BYTES >= count * THREADGROUP + ] + for count, size in cases: + measure(count, size) + + samples = {case: [] for case in cases} + for index in range(arguments.rounds): + ordered = cases[index % len(cases) :] + cases[: index % len(cases)] + for case in ordered: + samples[case].append(measure(*case)) + + rates = {case: statistics.median(values) for case, values in samples.items()} + + header = f"{'groups':>7}{'threads':>9}" + "".join( + f"{size // (1024 * 1024) if size >= 1024 * 1024 else size // 1024:>10}" + + ("MB" if size >= 1024 * 1024 else "KB") + for size in WORKING_SETS + ) + print(header) + print("-" * len(header)) + for count in THREADGROUP_COUNTS: + cells = [] + for size in WORKING_SETS: + rate = rates.get((count, size)) + cells.append(f"{rate:>10.0f} " if rate else f"{'-':>12}") + print(f"{count:>7}{count * THREADGROUP:>9}" + "".join(cells)) + + print("\nthreadgroups needed to reach 90% of the best this sweep saw, per working set:") + knees = {} + for size in WORKING_SETS: + curve = [ + (count, rates[count, size]) for count in THREADGROUP_COUNTS if (count, size) in rates + ] + peak = max(rate for _, rate in curve) + knee = next(count for count, rate in curve if rate >= 0.9 * peak) + knees[size] = knee + label = f"{size // (1024 * 1024)} MB" if size >= 1024 * 1024 else f"{size // 1024} KB" + print(f" {label:>7}: {knee:>4} groups ({knee * THREADGROUP} threads) for {peak:.0f} GB/s") + + # What this means for tile selection, which is the reason to measure it. + print("\nfor a GEMM tiled by output width, groups launched = output_features / block_n:") + print(f"{'N':>7}" + "".join(f"{f'bn={bn}':>10}" for bn in (32, 64, 128, 256))) + worst_knee = max(knees.values()) + for output_features in (4864, 8192, 8960, 17408): + cells = [] + for block_n in (32, 64, 128, 256): + groups = output_features // block_n + cells.append(f"{groups:>7}{' ok' if groups >= worst_knee else ' LOW'}") + print(f"{output_features:>7}" + "".join(cells)) + print(f"\nLOW marks fewer than {worst_knee} groups, the most demanding knee above.") + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/metile/target/__init__.py b/metile/target/__init__.py index 322d56d..c338703 100644 --- a/metile/target/__init__.py +++ b/metile/target/__init__.py @@ -7,6 +7,7 @@ from metile.target.agx import ( BANDWIDTH_BY_WORKING_SET_GBPS, + GROUPS_FOR_RESIDENT_BANDWIDTH, ILP_CEILING, MATRIX_PEAK_TFLOPS, REGISTER_BUDGET, @@ -31,6 +32,7 @@ __all__ = [ "BANDWIDTH_BY_WORKING_SET_GBPS", + "GROUPS_FOR_RESIDENT_BANDWIDTH", "ILP_CEILING", "MATRIX_PEAK_TFLOPS", "REGISTER_BUDGET", diff --git a/metile/target/agx.py b/metile/target/agx.py index 7bea96c..d1fd284 100644 --- a/metile/target/agx.py +++ b/metile/target/agx.py @@ -38,39 +38,43 @@ STREAMING_READ_GBPS = 120.6 # Read bandwidth as a function of working-set size, measured by benchmarks/agx_memory_hierarchy.py. -# One number for bandwidth is badly wrong here: a working set that stays resident is served sixteen to -# twenty times faster than one that streams, and that ratio is larger than every other factor in this -# file. A tiling that fits and a tiling that misses are not the same kernel. +# One number for bandwidth is badly wrong here: a working set the fast level holds is served twenty times +# faster than one that streams, and that ratio is larger than every other factor in this file. # -# Coalesced streaming reads, 16384 threads, six gigabytes of traffic per dispatch, sizes interleaved -# across rounds. The pass loop re-reads the same addresses, so the obvious worry is that the backend -# collapses it into a multiply and the fast numbers are fiction; tripling the traffic triples the -# elapsed time at every size, which it could not if the loop were collapsed. +# Each size is measured at several threadgroup counts and the best is kept, which is not a refinement but +# the difference between measuring the hierarchy and measuring the thread count. An earlier version used +# 64 threadgroups for every size and recorded a knee at 2 MB. There is no knee at 2 MB. At 64 groups an +# 8 MB working set reads 196 GB/s and at 512 it reads 2431, so what that table described was starvation: +# the thread count alone moves 16 MB by 19.7x. The fast level holds at least 16 MB. # -# The resident regime is one entry, not four, because the probe cannot resolve it. Measured 1545, 2006, -# 2138 and 2386 GB/s at 256 KB, 512 KB, 1 MB and 2 MB -- bandwidth *rising* with working set, which -# cannot be a property of a cache. It is the pass loop: a smaller working set means fewer inner -# iterations per pass, so loop bookkeeping takes a larger share, and the effect shrinks as the set -# grows. Only the 2 MB figure is close to uncontaminated, and it is a floor. -# -# So what this establishes is that a resident working set runs at 2386 GB/s or better, and nothing about -# how that varies below 2 MB. Reporting four numbers would claim a resolution the measurement does not -# have; a monotonicity test on the table is what caught the attempt. +# Sizes below about 1 MB are floors rather than levels. Enough threads to saturate and enough work per +# thread to amortise the pass loop are conflicting demands once the set is only a few times the thread +# count, so the resident regime is recorded as one figure. BANDWIDTH_BY_WORKING_SET_GBPS = { - 2 * 1024 * 1024: 2386.0, - 4 * 1024 * 1024: 555.0, - 8 * 1024 * 1024: 192.0, - 16 * 1024 * 1024: 161.0, - 32 * 1024 * 1024: 134.0, - 64 * 1024 * 1024: 128.0, - 128 * 1024 * 1024: 124.0, + 16 * 1024 * 1024: 2453.0, + 32 * 1024 * 1024: 1795.0, + 64 * 1024 * 1024: 448.0, + 128 * 1024 * 1024: 175.0, + 256 * 1024 * 1024: 156.0, } -# Largest working set still served by the fast level, and what it delivers. The knee is sharp: 2 MB -# reads 2386 GB/s and 4 MB reads 555, so this is the number a tiling pass should be trying to stay -# under. The rate is a floor for the reason above. -RESIDENT_WORKING_SET_BYTES = 2 * 1024 * 1024 -RESIDENT_READ_GBPS = 2386.0 +# Largest working set the fast level still serves, and what it delivers there. +RESIDENT_WORKING_SET_BYTES = 16 * 1024 * 1024 +RESIDENT_READ_GBPS = 2453.0 + +# Threadgroups needed before a resident working set reaches its bandwidth, by size. Below these the data +# is resident and the kernel still reads at streaming speed, which is the failure the corrected table +# above was hiding. +# +# Not yet shown to be actionable for real kernels, and worth saying so. The affine decode matmul launches +# `output_features / block_n` groups, which spans 35 to 280 at N=8960 -- an eightfold range straddling +# these thresholds -- and its configs measure within 3% of each other. Whatever binds that kernel, it is +# not this. The numbers describe streaming reads, where the effect is real and large. +GROUPS_FOR_RESIDENT_BANDWIDTH = { + 4 * 1024 * 1024: 256, + 8 * 1024 * 1024: 512, + 16 * 1024 * 1024: 1024, +} # Threadgroup memory, measured by benchmarks/agx_threadgroup_bandwidth.py against a resident device read # over the same 32 KB, so this compares staging with a cache hit rather than with DRAM. @@ -170,8 +174,9 @@ def tiling_gain(working_set_bytes): it. Neither of meTile's two main regimes has any: decode each weight element is read exactly once, so the working set is the whole weight and no - tiling changes that. Real MLP weights run 2.5 MB to 50 MB, all above the knee, and the - chosen configs achieve 80 to 128 GB/s against their footprint's level of 128 to 555. + tiling changes that. Real MLP weights run 2.5 MB to 50 MB, and while most of those are + inside the fast level, residency across decode steps does not survive: a token touches + every layer's weights, roughly a gigabyte, between two reads of the same one. prefill compute bound, not memory bound. The generated kernels reach 0.96x to 0.97x of MATRIX_PEAK_TFLOPS, and MLX reaches 0.95x to 0.96x, so there is nothing for a tiling to recover. diff --git a/tests/test_target.py b/tests/test_target.py index d7bd242..4295820 100644 --- a/tests/test_target.py +++ b/tests/test_target.py @@ -103,18 +103,18 @@ def test_the_backend_normalises_statement_order(): ) -def test_bandwidth_depends_on_working_set_and_the_knee_is_sharp(): +def test_bandwidth_depends_on_working_set_and_spans_more_than_ten_times(): """One bandwidth number is badly wrong for this part, so the model has to be a curve. - The measured drop from 2 MB to 4 MB is about a factor of four. Interpolating across it would invent - a smooth ramp the hardware does not have, so the lookup reports the measurement for the smallest - size at least as large as the request. + This used to assert the drop past the resident capacity was at least threefold, which encoded a knee + that measurement later removed. The first table was taken at a single thread count and put the + capacity at 2 MB with a fourfold cliff; sweeping occupancy per size moved the capacity to 16 MB and + made the far side a ramp -- 2453, 1795, 448, 175 -- so the sharp-knee assertion was testing an + artefact of the harness. What survives is the span and the ordering. """ from metile.target import RESIDENT_WORKING_SET_BYTES, read_bandwidth_gbps - resident_rate = read_bandwidth_gbps(RESIDENT_WORKING_SET_BYTES) - beyond_rate = read_bandwidth_gbps(RESIDENT_WORKING_SET_BYTES * 2) - assert resident_rate > 3 * beyond_rate + assert read_bandwidth_gbps(RESIDENT_WORKING_SET_BYTES) > 10 * STREAMING_READ_GBPS assert read_bandwidth_gbps(2**30) == pytest.approx(STREAMING_READ_GBPS) # Monotonically non-increasing: a larger working set is never served faster. This is what caught @@ -125,6 +125,19 @@ def test_bandwidth_depends_on_working_set_and_the_knee_is_sharp(): assert faster >= slower +def test_reaching_resident_bandwidth_takes_parallelism_as_well_as_residency(): + """Residency is necessary and not sufficient, which is what the first table missed. + + At 64 threadgroups an 8 MB working set reads 196 GB/s; at 512 it reads 2431. Data being resident buys + nothing if too few threadgroups are in flight to ask for it. + """ + from metile.target import GROUPS_FOR_RESIDENT_BANDWIDTH + + required = sorted(GROUPS_FOR_RESIDENT_BANDWIDTH.items()) + assert [groups for _, groups in required] == sorted(groups for _, groups in required) + assert all(groups >= 256 for _, groups in required) + + def test_a_working_set_is_resident_only_up_to_the_measured_capacity(): from metile.target import RESIDENT_WORKING_SET_BYTES, resident