diff --git a/dm/games/0400486024DEC000/README.md b/dm/games/0400486024DEC000/README.md
index d269496e33..3513718f6e 100644
--- a/dm/games/0400486024DEC000/README.md
+++ b/dm/games/0400486024DEC000/README.md
@@ -74,6 +74,7 @@ From an empty save: at about 45 s `A` (title), `A` (New Game), then the movie. F
 
 ## Found on this game
 
+- [Raw global memory lookups that reuse their last range](../../../docs/reference/shaders.md#raw-global-memory-lookups-that-reuse-their-last-range): the foliage, vertex and compute shaders read structures word by word through loaded addresses, each word with its own lookup of the GMEM map.
 - [The GPU idles with the next command buffer already committed](../../../docs/reference/frame-rate.md#the-gpu-idles-with-the-next-command-buffer-already-committed): the 25-40 ms frames.
 - [Pool tables and readback buffers made anew each frame](../../../docs/reference/memory.md#pool-tables-and-readback-buffers-made-anew-each-frame): about 1,500 residency changes per 300 frames.
 - [Pool shaders taken as depth readers](../../../docs/reference/rendering.md#pool-shaders-taken-as-depth-readers): about 80 extra passes a frame, and half the frames over budget.
diff --git a/dm/gfx/nvn/dmsh_chain.h b/dm/gfx/nvn/dmsh_chain.h
index 34b06ab0dd..1d067570da 100644
--- a/dm/gfx/nvn/dmsh_chain.h
+++ b/dm/gfx/nvn/dmsh_chain.h
@@ -1,3 +1,3 @@
 // Generated file (dm/gfx/tools/dmsh_chain.py --header): the B shader library translation chain version this source tree expects (default options), one for every game;
 // the host loads only <library dir>/b-<V>/. Run dmsh_chain.py --header after changing the chain's tools and commit this file with the change
-#define DMSH_CHAIN "adec7d81d970"
+#define DMSH_CHAIN "15a7cccf7a6b"
diff --git a/dm/gfx/tests/test_dmsh_raw_ranges.py b/dm/gfx/tests/test_dmsh_raw_ranges.py
new file mode 100644
index 0000000000..9eb937f9f2
--- /dev/null
+++ b/dm/gfx/tests/test_dmsh_raw_ranges.py
@@ -0,0 +1,122 @@
+"""Raw global memory lookups that reuse the range they last chose (dmsh_wide.py raw_ranges): the helper's search behind the cache, each call site's cache, shapes left
+alone, and the compile."""
+from pathlib import Path
+import os
+import shutil
+import subprocess
+import sys
+import tempfile
+import unittest
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / 'tools'))
+from dmsh_wide import raw_ranges   # noqa: E402
+
+# GlobalToStorage.TryGenerateRawStorageOp's load helper as spirv-cross writes it, its unrolled search cut to two steps, and a kernel that reads two neighbouring words
+KERNEL = '''#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _27
+{
+    uint _m0[1];
+};
+
+struct spvDescriptorSetBuffer0
+{
+    const device _27* cp_s3 [[id(0)]];
+    device _27* cp_s4 [[id(1)]];
+};
+
+static inline __attribute__((always_inline))
+int _56(thread const int& _100, thread const int& _101, const device _27& cp_s3)
+{
+    int _102 = _101 << 11;
+    uint _103 = uint(int(uint(_100) >> uint(21)));
+    int _104 = _102 | int(_103);
+    int _105 = _104 & 8191;
+    int _106 = _105 * 8;
+    int _107 = _106 + 6;
+    float _108 = as_type<float>(cp_s3._m0[_107]);
+    int _109 = as_type<int>(_108) & 8191;
+    bool _110 = _109 != 0;
+    int _111 = _109;
+    if (!_110)
+    {
+        int _112 = 0 + 4096;
+        int _113 = _112 * 8;
+        float _114 = as_type<float>(cp_s3._m0[_113]);
+        bool _115 = as_type<uint>(_114) <= uint(_100);
+        int _116 = _115 ? _112 : 0;
+        int _117 = _116 + 2048;
+        int _118 = _117 * 8;
+        float _119 = as_type<float>(cp_s3._m0[_118]);
+        bool _120 = as_type<uint>(_119) <= uint(_100);
+        int _121 = _120 ? _117 : _116;
+        _111 = _121;
+    }
+    int _122 = _111 * 8;
+    int _123 = _122 + 2;
+    float _124 = as_type<float>(cp_s3._m0[_123]);
+    bool _125 = uint(_100) < as_type<uint>(_124);
+    if (_125)
+    {
+        int _126 = _122 + 4;
+        float _127 = as_type<float>(cp_s3._m0[_126]);
+        int _128 = _100 + as_type<int>(_127);
+        uint _129 = *(reinterpret_cast<const device uint*>(as_type<ulong>(uint2(uint(_128), uint(_101)))));
+        float _130 = as_type<float>(_129);
+        return as_type<int>(_130);
+    }
+    return 0;
+}
+
+kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    int _200 = int(gl_GlobalInvocationID.x) * 8;
+    int _201 = 0;
+    int _202 = _56(_200, _201, (*spvDescriptorSet0.cp_s3));
+    int _203 = _200 + 4;
+    int _204 = _56(_203, _201, (*spvDescriptorSet0.cp_s3));
+    (*spvDescriptorSet0.cp_s4)._m0[gl_GlobalInvocationID.x] = uint(_202 + _204);
+}
+'''
+
+
+class RawRanges(unittest.TestCase):
+    def test_search_behind_the_cache(self):
+        out = raw_ranges(KERNEL)
+        self.assertIn('struct dm_raw_range { ulong s; ulong e; int i; };', out)
+        self.assertIn('int _56(thread const int& _100, thread const int& _101, const device _27& cp_s3, thread dm_raw_range& _rr)', out)
+        helper = out[out.index('int _56('):out.index('kernel void main0')]
+        # the hint and the search run only on a miss, and a miss records the chosen entry's range
+        self.assertLess(helper.index('if (_ra - _rr.s < _rr.e - _rr.s) _111 = _rr.i;'), helper.index('float _108 = '))
+        self.assertIn('        _111 = _109;', helper)
+        self.assertIn('_rr.i = _111;', helper)
+        # what follows the choice is unchanged
+        self.assertIn('    int _122 = _111 * 8;\n    int _123 = _122 + 2;', helper)
+
+    def test_call_sites_share_their_function_cache(self):
+        out = raw_ranges(KERNEL)
+        main = out[out.index('kernel void main0'):]
+        self.assertEqual(main.count('dm_raw_range _rr_56 = {0ul, 0ul, 0};'), 1)
+        self.assertIn('_56(_200, _201, (*spvDescriptorSet0.cp_s3), _rr_56)', main)
+        self.assertIn('_56(_203, _201, (*spvDescriptorSet0.cp_s3), _rr_56)', main)
+
+    def test_left_alone(self):
+        self.assertIs(raw_ranges('kernel void main0() {}'), 'kernel void main0() {}')
+        other = KERNEL.replace('    if (!_110)\n', '    if (_110)\n')   # not the helper's shape: the shader stays as it is
+        self.assertIs(raw_ranges(other), other)
+
+    @unittest.skipUnless(shutil.which('xcrun') and sys.platform == 'darwin', 'Metal compiler')
+    def test_compiles(self):
+        with tempfile.TemporaryDirectory() as d:
+            m = os.path.join(d, 'k.metal')
+            open(m, 'w').write(raw_ranges(KERNEL))
+            r = subprocess.run(['xcrun', '-sdk', 'macosx', 'metal', '-c', '-O1', '-fpreserve-invariance', m, '-o', os.path.join(d, 'k.air')],
+                               capture_output=True, text=True)
+            self.assertEqual(r.returncode, 0, r.stderr[-2000:])
+
+
+if __name__ == '__main__':
+    unittest.main()
diff --git a/dm/gfx/tools/dmsh_build.py b/dm/gfx/tools/dmsh_build.py
index a41265d20c..694dedabcc 100644
--- a/dm/gfx/tools/dmsh_build.py
+++ b/dm/gfx/tools/dmsh_build.py
@@ -1480,6 +1480,8 @@ def finish(key, w, j, c1):
     fu, fmask = fs_uint(src) if j['stage'] == 'ps' else (None, 0)   # the integer-output entry points (see fs_uint); the analysis below uses the source with only main0
     fd = fs_dual(src) if j['stage'] == 'ps' and key in DUAL_WANT else None   # the dual-source entry point (see fs_dual), from the source with only main0
     if fu or fd: plain = src; open(msl, 'w').write((fu or src) + (fd or ''))
+    final = open(msl).read(); rr = dmsh_wide.raw_ranges(final)   # raw global memory lookups reuse the range they last chose (see dmsh_wide.py raw_ranges)
+    if rr is not final: open(msl, 'w').write(rr)
     side, lost = {}, []
     def compile_targets():   # the selected targets from the same MSL (MSL identical to one compiled before takes that library, dmsh_mtlcache.py)
         if j['stage'] == 'ps': side.clear(); side.update(ps_libraries(key, msl, open(msl).read())); return
@@ -1489,7 +1491,7 @@ def compile_targets():   # the selected targets from the same MSL (MSL identical
     except subprocess.CalledProcessError as e:
         if not lay and not fu and not fd: return 'metal failed', str(e)[:160], None
         lost = [name for name, extra in (('layer', lay), ('integer output', fu), ('dual source', fd)) if extra]   # kept in the reflection, as a cache hit compiles nothing: every build reports it (shader_regressions.report_losses)
-        open(msl, 'w').write(plain); lay = fu = fd = None   # the layer, integer-output and dual-source entry points are extras: a library that does not build with them is built without
+        open(msl, 'w').write(dmsh_wide.raw_ranges(plain)); lay = fu = fd = None   # the layer, integer-output and dual-source entry points are extras: a library that does not build with them is built without
         try: compile_targets()
         except subprocess.CalledProcessError as e: return 'metal failed', str(e)[:160], None
     clone(msl, os.path.join(out, key + '_b.metal'))
diff --git a/dm/gfx/tools/dmsh_wide.py b/dm/gfx/tools/dmsh_wide.py
index 971af9db8c..f73620cc8c 100644
--- a/dm/gfx/tools/dmsh_wide.py
+++ b/dm/gfx/tools/dmsh_wide.py
@@ -73,3 +73,78 @@ def wide_loads(msl):
         indent = lines[i][:len(lines[i]) - len(lines[i].lstrip())]
         lines[i] = ''.join(indent + d + '\n' for d in pre) + ln
     return '\n'.join(lines)
+
+
+# Raw global memory that reuses the range its last lookup found (raw_ranges, the last step before dmsh_build.py compiles a shader).
+#
+# A word whose address the translator traced to no storage buffer goes through a helper GlobalToStorage.TryGenerateRawStorageOp writes (spirv-cross keeps it a function):
+#
+#     static inline __attribute__((always_inline))
+#     int _56(thread const int& lo, thread const int& hi, const device _27& cp_s3)
+#     {
+#         ... the hint of the address's 2 MiB chunk ...      int sel = hint_index;
+#         if (!hint_fits) { ... the unrolled binary search ...  sel = index; }
+#         int row = sel * 8;  ... the end check, the delta, the read ...
+#     }
+#
+# called once for each 32-bit word, and a shader reads a structure word after word from one range. Ocarina of Time's foliage mesh shader (73dc0b90c4082e43, 135 such reads)
+# took 10.9 ms of a 683x384 frame. Each call site gets the cache of the function it is in, a range {start, end, entry}: an address inside the range the last lookup chose takes
+# that entry without the hint or the search; everything after (the end check, the delta, the read or write) is unchanged. The GMEM map's ranges are disjoint (nvn_mtl.m gmem
+# map), so an address inside an entry's range is that entry's in the search too, and a cache lives in one invocation, so the table cannot change under it.
+_RAW_HELPER = re.compile(r'(?m)^static inline __attribute__\(\(always_inline\)\)\nint (_\d+)\(thread const int& (_\d+), thread const int& (_\d+), '
+                         r'((?:const )?device _\d+& ([a-z]p_s\d+))\)\n\{\n')
+_RAW_CACHE = 'struct dm_raw_range { ulong s; ulong e; int i; };   // the range the last raw global memory lookup chose (dmsh_wide.py raw_ranges)\n\n'
+
+
+def _raw_helper(msl, m):
+    """The helper with its search behind the cache, and where it ends; None when it is not the shape above."""
+    name, lo, hi, param, table = m.groups()
+    end = msl.index('\n}\n', m.end())
+    body = msl[m.end():end].split('\n')
+    i = next((k for k in range(len(body) - 2) if re.fullmatch(r'    int _\d+ = _\d+;', body[k]) and re.fullmatch(r'    if \(!_\d+\)', body[k + 1])
+              and body[k + 2] == '    {'), None)
+    if i is None: return None
+    sel, hint = re.fullmatch(r'    int (_\d+) = (_\d+);', body[i]).groups()
+    close = next((k for k in range(i + 3, len(body)) if body[k] == '    }'), None)
+    if close is None or close + 1 >= len(body) or not re.fullmatch(r'    int _\d+ = ' + sel + r' \* 8;', body[close + 1]): return None
+    new = ([f'    ulong _ra = as_type<ulong>(uint2(uint({lo}), uint({hi})));', f'    int {sel};',
+            f'    if (_ra - _rr.s < _rr.e - _rr.s) {sel} = _rr.i;', '    else', '    {']
+           + ['    ' + ln for ln in body[:i]] + [f'        {sel} = {hint};'] + ['    ' + ln for ln in body[i + 1:close + 1]]
+           + [f'        int _rw = {sel} * 8;',
+              f'        _rr.s = as_type<ulong>(uint2({table}._m0[_rw], {table}._m0[_rw + 1])); _rr.e = as_type<ulong>(uint2({table}._m0[_rw + 2], {table}._m0[_rw + 3])); _rr.i = {sel};',
+              '    }'] + body[close + 1:])
+    head = msl[m.start():m.end()].replace(f'{param})\n{{\n', f'{param}, thread dm_raw_range& _rr)\n{{\n')
+    return name, head + '\n'.join(new), end
+
+
+def _close_paren(s, i):
+    depth = 0
+    for j in range(i, len(s)):
+        depth += (s[j] == '(') - (s[j] == ')')
+        if depth == 0: return j
+    raise ValueError('unbalanced parentheses')
+
+
+def raw_ranges(msl):
+    if 'reinterpret_cast<' not in msl: return msl
+    helpers, out, pos = [], [], 0
+    for m in _RAW_HELPER.finditer(msl):
+        r = _raw_helper(msl, m)
+        if not r: return msl   # a shape this does not know: the shader stays as it is
+        name, text, end = r
+        out += [msl[pos:m.start()], text]; pos = end
+        helpers.append(name)
+    if not helpers: return msl
+    src = ''.join(out) + msl[pos:]
+    opens = [m.start() for m in re.finditer(r'(?m)^\{$', src)]   # function bodies open at column 0
+    edits, decl = [], {}
+    for name in helpers:
+        for m in re.finditer(r'(?<![\w.])' + name + r'\(', src):
+            if src.startswith('int ' + name + '(', src.rfind('\n', 0, m.start()) + 1): continue   # its definition
+            edits.append((_close_paren(src, m.end() - 1), f', _rr{name}'))
+            decl.setdefault(max(p for p in opens if p < m.start()), set()).add(name)
+    if not edits: return msl
+    for at, names in decl.items(): edits.append((at + 2, ''.join(f'    dm_raw_range _rr{n} = {{0ul, 0ul, 0}};\n' for n in sorted(names))))
+    for at, text in sorted(edits, key=lambda e: -e[0]): src = src[:at] + text + src[at:]
+    first = src.index('static inline __attribute__((always_inline))\nint ' + helpers[0] + '(')
+    return src[:first] + _RAW_CACHE + src[first:]
diff --git a/docs/agents/symptoms.md b/docs/agents/symptoms.md
index a308c2dd59..964eb59bcb 100644
--- a/docs/agents/symptoms.md
+++ b/docs/agents/symptoms.md
@@ -879,6 +879,7 @@ One line per symptom in the [subsystem references](../reference/): what you see
 - bindless handles read from a uniform buffer name the wrong texture when the game writes them earlier in the same submission. → [Handles in uniform buffers written earlier in the same submission](../reference/shaders.md#handles-in-uniform-buffers-written-earlier-in-the-same-submission)
 - a storage buffer reported with `cb_slot` 3, which had no NVN slot. → [Storage buffer addresses in the game's own uniform buffer](../reference/shaders.md#storage-buffer-addresses-in-the-games-own-uniform-buffer)
 - whatever a shader writes through a pointer it read from memory never happens. The build summary counted `Failed to find storage buffer for global memory operation "Store"`, and the shader's MSL had a store function with an empty body. In Donkey Kong Bananza the tunnel's walls and floor were black and only the loose rocks drew. → [Global memory at an address loaded from memory](../reference/shaders.md#global-memory-at-an-address-loaded-from-memory)
+- a pass whose shaders read memory through loaded addresses takes several milliseconds on the GPU for little work. `[gpuprof]` names the pass, and its shaders' build notes say `Global memory operation "Load" traced to no storage buffer`. Ocarina of Time's foliage mesh shader (`73dc0b90c4082e43`, `Foliage.Product.110`, 135 such reads) took 10.9 ms of a 683x384 frame. 2,228 of the game's 42,853 shaders have the helper (vertex, compute, fragment and mesh), and some call it 971 times. → [Raw global memory lookups that reuse their last range](../reference/shaders.md#raw-global-memory-lookups-that-reuse-their-last-range)
 - `no resolved slot number`, the shader does not load; the translation has a 32-entry texture-pool array (`fp_t2d_cb0_40000000`, handle type Direct). → [Bindless handles through a loop phi](../reference/shaders.md#bindless-handles-through-a-loop-phi)
 - none in the log: texture reads come back 0. Particle meshes collapse to a point (their scale reads 0), so effects are missing or show as dark patches. The translator's own message (`Failed to find handle source for bindless access of type "textureBuffer"`) used to be dropped, because `dmsh_xlate`'s accessor had no `Log`; its messages now go into the shader's notes, marked `translator: `, and the build summary counts them apart from unanswered queries. → [Bindless reads whose handle phi also carries other values](../reference/shaders.md#bindless-reads-whose-handle-phi-also-carries-other-values)
 - the shaders' bindless arrays have no slot and used not to load; the translator falls back to a 32-entry array such as `array<texture_buffer<float>, 32> vp_tb_cb0_40000000`. → [Texture slots picked at run time](../reference/shaders.md#texture-slots-picked-at-run-time)
diff --git a/docs/reference/shaders.md b/docs/reference/shaders.md
index e1ca20af2c..5555e3e8fd 100644
--- a/docs/reference/shaders.md
+++ b/docs/reference/shaders.md
@@ -1399,6 +1399,14 @@ Finding, keying and translating shaders (`dm/gfx/tools/shader_split.py`, `dmsh_b
 - **Open:** the prediction is checked, not proven: a miss that does reach a shadow is reported (`[RED]`) after the frame, not prevented. The records' pointers come from memory, so the pools they will reach are not known ahead of the draw.
 - **Found on:** [Donkey Kong Bananza](../../dm/games/040074A01BF12000/README.md): its terrain, and agl's compute shader `49bc2268925e25b6` (one load); rebuilding the chain changed only these two of Donkey Kong Bananza's shaders. Vertex shader `588e242a52f5f7aa` (five stores) has one too, the same shader in [The Legend of Zelda: Tears of the Kingdom](../../dm/games/0100F2C0115B6000/README.md) and [Super Mario Bros. Wonder](../../dm/games/010015100B514000/README.md), so likely one the SDK ships; neither drew it in its regression route.
 
+### Raw global memory lookups that reuse their last range
+
+- **Symptom:** a pass whose shaders read memory through loaded addresses takes several milliseconds on the GPU for little work. `[gpuprof]` names the pass, and its shaders' build notes say `Global memory operation "Load" traced to no storage buffer`. Ocarina of Time's foliage mesh shader (`73dc0b90c4082e43`, `Foliage.Product.110`, 135 such reads) took 10.9 ms of a 683x384 frame. 2,228 of the game's 42,853 shaders have the helper (vertex, compute, fragment and mesh), and some call it 971 times.
+- **Cause:** each 32-bit word goes through `GlobalToStorage.TryGenerateRawStorageOp`'s helper on its own: the hint of its 2 MiB chunk, and on a miss the unrolled 13-step search of the GMEM map. A shader reads a structure word after word from one range, so all but the first lookup choose the entry the first one chose.
+- **Fix:** `dmsh_wide.py` `raw_ranges`, the last step before `dmsh_build.py` compiles a shader. Each function that calls the helper keeps the range its last lookup chose (`dm_raw_range`: start, end, entry), and an address inside it takes that entry without the hint or the search. Everything after the choice (the end check, the delta, the read or write) is unchanged. The GMEM map's ranges are disjoint, so the search would choose the same entry, and the cache lives in one invocation, so the table cannot change under it. A helper of another shape leaves the shader as it is.
+- **Check:** `dm/gfx/tests/test_dmsh_raw_ranges.py`. Ocarina of Time on a MacBook Pro M5 Pro (chain `adec7d81d970`, the 2,228 libraries replaced by loose ones, the same scripted seven-minute route into Kokiri Forest, the pipeline archives warmed by a run before): GPU per frame p50 17.0 → 14.1 ms, frame time p50 / p95 17.3 / 18.4 → 16.5 / 18.1 ms, frames dropped on screen 9 → 0 per 300; one frame's encoders 17.5 → 15.0 ms from first start to last end. Screenshots at the spawn point match.
+- **Found on:** [The Legend of Zelda: Ocarina of Time](../../dm/games/0400486024DEC000/README.md).
+
 ### Bindless handles through a loop phi
 
 - **Symptom:** `no resolved slot number`, the shader does not load; the translation has a 32-entry texture-pool array (`fp_t2d_cb0_40000000`, handle type Direct).