From dfc720e7547fd6f056bed4a589c73d6088dc631c Mon Sep 17 00:00:00 2001 From: Meta Vens Date: Wed, 14 Oct 2026 06:05:27 +0200 Subject: [PATCH] Research: tape register model prior art (Fidget, MPR, GPU indexed locals) and body-Voronoi metrics Eyes pass (2026-10-14) serving cultmath-tapes:question:tape-register-model, with a scope addition on weighted diagrams over gravity bodies and metric parameterisation for question lattice-site-weight-source. Co-Authored-By: Claude Opus 5.5 --- docs/research/tape-register-model.md | 82 +++++++++++++++++++++ docs/research/tape-register-model/probe2.py | 52 +++++++++++++ 2 files changed, 134 insertions(+) create mode 100644 docs/research/tape-register-model.md create mode 100644 docs/research/tape-register-model/probe2.py diff --git a/docs/research/tape-register-model.md b/docs/research/tape-register-model.md new file mode 100644 index 00000000..3538245e --- /dev/null +++ b/docs/research/tape-register-model.md @@ -0,0 +1,82 @@ +# Facts for `cultmath-tapes:question:tape-register-model` (2026-10-14) + +Eyes, read only. Tags: [src] source read, [doc] vendor/library documentation, [web] web page or search summary (second-hand where noted), [abs] derived by reasoning or a synthetic probe, not a measurement of any real GPU. Raw downloads: `raw/` beside this file (Fidget files at `main` on 2026-10-14; MPR paper text). Probe: `probe2.py` beside this file. + +## 1. Fidget and MPR + +### 1a. Fidget tape, allocator, bytecode (`github.com/mkeeter/fidget`, raw at `raw.githubusercontent.com/mkeeter/fidget/main/`) +- [src] The allocator is `fidget-core/src/compiler/alloc.rs`, `RegisterAllocator`, doc "Cheap and cheerful single-pass register allocation" (`:11`). `assert!(N <= u8::MAX)` (`:51`). It keeps a least-recently-used list of registers (`register_lru`) and the output tape is "assembled in reverse order". When no register is free it spills to a memory slot (`Allocation::Memory(u32)`, `RegOp::Load/Store`); memory is "treated as unlimited" (`get_memory`, `:106-123`). So the allocator is not a cap-and-refuse design: the CPU/JIT target spills past N. (I did not read the walk direction beyond the "reverse order" comment; "linear scan from the tape end" is consistent with it, not confirmed line by line.) +- [src] Bytecode, `fidget-bytecode/src/lib.rs:1-45`: little-endian `u32` words, two words per op (fixed length so it can be walked forwards and backwards). Word 0 bytes: opcode, output register, input register 1, input register 2; input byte `0xFF` means "use word 1 as an f32 immediate". Tape starts `0xFFFFFFFF 0x00000000`, ends `0xFFFFFFFF 0xFFFFFFFF`. `Mem` op (Load/Store) uses word 1 as a memory slot. Register 255 is reserved (`ReservedRegister` error), so at most 255 registers in bytecode. The format is stated "not stable". Registers are renumbered by frequency of use (most used = 0). `Bytecode` records `reg_count` and `mem_count`. +- [src] `fidget-wgpu/src/shaders/tape_interpreter.wgsl` (115 lines): `fn run_tape(...)` declares `var reg: array;` as a function-scope (private, per-invocation) array, indexed at runtime by bytes from the op word (`reg[lhs_i]`, `reg[op[1]] = tmp`). Not workgroup memory. `REG_COUNT` is a WGSL `const` the host prepends per pipeline (`fidget-wgpu/src/pixel/mod.rs:37,54,71`: `const REG_COUNT: u32 = {reg_count};`). +- [src] `fidget-wgpu/src/lib.rs:529-545`: `REG_PIPELINE_SIZES = [8, 16, 32, 64, 128, 192, 255]`; `RegPipeline::get(reg_count)` compiles (lazily, cached) the first pipeline whose size >= the tape's register count and panics above 255. So the GPU interpreter specialises the shader on register-array length in seven buckets, so small tapes do not pay for a 255-element array. +- [src] The same interpreter does NOT support spills: `case OP_MEM: { // Not implemented! return out; }` (op switch in `tape_interpreter.wgsl`). So on the GPU the effective cap is 255 registers and a tape needing memory slots does not run correctly. Related capacity: `TAPE_DATA_CAPACITY` = 8M words, 64 MiB (`lib.rs:268`). +- [src] Registers hold `Value` (shared across interval, float and gradient evaluation). Pixel workgroups are `@workgroup_size(8, 8)` (`pixel/shaders/pixel_tiles.wgsl:7`, `interval_tiles.wgsl:13`). No `var` register store in those two shaders (grep found none). +- [src] A second per-thread private structure: the min/max choice `Stack` is `array` per invocation, 2 bits per choice, 512 choices (`shaders/stack.wgsl:1-12`). +- [web] I found no published performance numbers for `fidget-wgpu` (its README is 9 lines: "implements 3D rendering of Fidget tapes using WebGPU"). CultLib's own note: `F:\Projects\CultLib\docs\cultmath-tape-target.md:25-40` (Fidget 0.5.1; tape 254 clauses shrinks to about 20 at pixel level, Fidget's example). + +### 1b. MPR (Keeter, ACM TOG 39(4), SIGGRAPH 2020; `https://www.mattkeeter.com/research/mpr/keeter_mpr20.pdf`, text in `raw/mpr.txt`) +- [src] Section 2: each thread "has exclusive access to a block of memory"; clause arguments and outputs are indexes into it, "slots", "equivalent to machine registers". Slot allocation is "a simple greedy algorithm": a slot is free outside its value's live range; if an argument is the last use, its slot is freed before the output is assigned (a clause may read and write the same slot). Register allocation with reordering (Rawat et al. 2018) was "deemed outside the scope". +- [src] Section 4 implementation list: clause = 8 bytes (opcode 1 B, output slot 1 B, two input slots 1 B each, or one slot + 4 B immediate); "The storage slots are per-thread arrays of 128 values. For simplicity, this is hard-coded, but could be dynamically sized with the addition of load/store opcodes." The choices stack is a per-thread array too (2 bits per choice, 4096 choices in 256 `uint32_t`). Implementation was CUDA; tape in global memory. +- [src] Table 2 (clauses, slots, ratio): "Hello, world" 328 / 26 / 12.6; bear head 541 / 22 / 24.6; architectural model 961 / 75 / 12.8; 2D text benchmark 6056 / 84 / 72.1. The paper hypothesises the high text ratio is because each character is an independent CSG primitive. +- [src] Interpreter overhead (section 2.1, GTX 1080 Ti): brute-force interpreter about 58 ms per megapixel versus about 3 ms/MP for the same expression pasted into a CUDA function and compiled with nvcc (about 19x), on the 6056-clause text model; nvcc compile took 3.45 s. The full algorithm (interval pruning, tape shortening) beats the compiled kernel above about 1536^2 pixels and stays below 8 ms/frame at 4096^2. The paper reports no separate measurement of the 128-slot array's local-memory cost. +- [src] Tape generation on CPU: 7 ms for the text model, 1-2 ms for the others. + +## 2. GPU cost of dynamically indexed locals, groupshared, unrolled registers + +- [doc] NVIDIA CUDA Programming Guide (`https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html`, via search summary): the compiler is likely to place arrays in local memory when it cannot prove constant indices; local memory physically lives in global memory space. The Nsight Compute profiling guide says the same. +- [web] NVIDIA developer blog "Fast Dynamic Indexing of Private Arrays in CUDA" (`https://developer.nvidia.com/blog/fast-dynamic-indexing-private-arrays-cuda/`, fetched): constant indices keep the array in registers; dynamic indexing forces local memory, cached in L1/L2. Measured on Tesla K20, 32-element float private array: non-uniform indexing in local memory 24x slower than uniform indexing in local memory; non-uniform indexing in shared memory 1.34x faster than uniform local. On Maxwell: local non-uniform vs uniform gap about 14x; the shared-memory version 4x faster than uniform local. Replays per warp: 0 if all lanes use the same index, about 2.5 average for uniformly random indices over 0..31, up to 31 if every lane differs. Suggests a math-to-memory ratio of roughly 4:1 to 8:1 to hide local-memory cost. Shared layout: element i of thread t at `i * BLOCK + t` (one bank per thread); block size a multiple of 32; at most 48 KB shared per block; technique suited to private arrays of about 30-50 32-bit elements, larger arrays reduce occupancy. +- [web] AMD (search summary of ROCm docs and a 2012 GCN thread): stack arrays live in scratch memory by default and the compiler may promote them to registers; GCN VGPRs can be indexed but one contributor reported about 50% speedup moving to static registers in some kernels (anecdote, old arch: `https://www.openwall.com/lists/john-dev/2012/12/09/7`). LLVM has a flag `-amdgpu-use-divergent-register-indexing`, so divergent VGPR indexing is a compiler mode. Register budget vs occupancy: RDNA SIMD register file 128 KB (GCN 64 KB); RDNA 4 desktop 192 KB per SIMD, full thread slots at 96 or fewer vector registers (`https://chipsandcheese.com/p/dynamic-register-allocation-on-amds`, `https://rocm.blogs.amd.com/software-tools-optimization/register-pressure/README.html`). +- [web] Intel oneAPI GPU optimization guide, "Optimizing register spills" (`https://www.intel.com/content/www/us/en/docs/oneapi/optimization-guide-gpu/2024-2/optimizing-register-spills.html`): spills are likelier on Intel than NVIDIA because the GRF is smaller (A100 up to 256 32-bit registers per work-item vs 64 by default on Data Center GPU Max 1550); advice to avoid big private arrays and use shared local memory. A 2013 Intel SDK note says indexed private memory is slow because each work-item touches a unique cache line (old, OpenCL). +- [web] Mesa/NIR driver commits (search summary): drivers choose by a size threshold between lowering indirect array access to conditional register selects and lowering to scratch. Intel commit "intel/fs: Lower large local arrays to scratch": on Kaby Lake, helped 6, hurt 2, helped programs were Aztec Ruins compute shaders, overall confidence interval for cycles included zero. An AMD MR found a 16-scalar-element threshold caused more scratch use. V3DV keeps arrays larger than 2 elements in scratch. SPIR-V allows dynamic indexing of arrays in every storage class except PushConstant (Khronos slide deck). Links: `https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/1097`, `https://gitlab.neroreflex.duckdns.org/NeroReflex/mesa/commit/69244fc7`. +- [doc] D3D: `dcl_indexableTemp x0[23], 2` declares an indexable temp array; element count of all temp registers (indexable and not) cannot exceed 4096 (`https://learn.microsoft.com/en-us/windows/win32/direct3dhlsl/dcl-indexabletemp`). FXC names a locally indexed array `x#`, each its own namespace (`https://list.winehq.org/hyperkitty/list/wine-gitlab@list.winehq.org/message/KNVE3CV4O2D464TKNCRIMZA6B7QG54NP/`). I found no measured cost of `x#` access on any GPU; whether a driver puts an `x#` array in registers or scratch is the driver's choice (vendor behaviour above). +- [web] FXC limitation: writing to a dynamically indexed fixed-size array inside a struct fails (X3500 "array reference cannot be used as an l-value; not natively addressable"); Tint copies the array to a local, writes, copies back (`https://dawn.googlesource.com/tint/+/refs/heads/chromium/4765%5E%21/`). Relevant if the register array sits in a struct. +- [doc] HLSL `groupshared`: compute shaders only; D3D10 limit 16 KB, D3D11 limit 32 KB total per thread group (`https://learn.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-variable-syntax`). [abs] 32 KB / (N registers x 4 B) = threads per group: N=64 gives 128 threads; N=128 gives 64; N=255 gives 32 (a float2 interval register halves each). +- [abs] Static indexing needs the register number to be a compile-time constant, which in an interpreter loop it is not. A fully unrolled generated kernel has constant registers; that is MPR's "compiled expression" case (3 ms/MP vs 58 ms/MP interpreted, 1080 Ti). + +## 3. Other GPU interpreters and expression VMs + +- [web] Unity Shader Graph: the node graph is emitted as HLSL; the Custom Function node injects HLSL text or an include (`https://docs.unity3d.com/Packages/com.unity.shadergraph@17.2/manual/Custom-Function-Node.html`). The docs I read do not describe a runtime interpreter; the internals of the generator are not stated there, so "no interpreter" is consistent with the docs, not confirmed by them. Built-in pipeline support in Shader Graph is "compatibility only" since 2021.2 and, in the 17.5 docs, deprecated, supported through the Unity 6.7 LTS lifecycle (`https://docs.unity3d.com/Packages/com.unity.shadergraph@17.5/manual/Graph-Target.html`). Shader Graph docs found say nothing about compute shaders. +- [web] Dreams (Media Molecule): Alex Evans, SIGGRAPH 2015 "Learning from Failure" (`https://advances.realtimerendering.com/s2015/AlexEvans_SIGGRAPH-2015-sml.pdf`, `https://www.mediamolecule.com/blog/article/siggraph_2015`). Secondary summary: edit lists of SDF primitives evaluated hierarchically into cached 8x8x8 bricks. The slides exceeded the fetch size cap (10 MB), so not read; no source found for a bytecode interpreter or its register storage in Dreams. GAP. +- [web] vyre crate (`https://docs.rs/crate/vyre/0.3.0`): compiles boolean rule conditions to bytecode (49 opcodes) run on compute shaders, one rule per thread; no register layout described. +- [src] Per `F:\Projects\CultLib\docs\cultmath-interval-ground-cut.md:1727-1731`, `F:\Projects\aetheria-material-path-prior-art.md` section 4 holds Dolphin ubershaders and Laine et al. on megakernel register pressure; "no measured interpreter-versus-kernel timing was found". I did not re-read that file. +- [web] Divergence: an interpreter loop dispatches on data-dependent opcodes; Mesa ACO notes divergent control flow executes both paths with lanes masked (`https://android.googlesource.com/platform/external/mesa3d/+/b2fa2d99ae2/src/amd/compiler/README.md`). MPR makes every thread of a warp evaluate the same shortened tape (MPR section 4). [abs] A fixed trophic graph evaluated per point has the same property: all lanes run the same tape. +- Gaps: no hit for libfive GPU work other than MPR itself; no IQ note on bytecode interpreters; no published shadertoy-style VM register layout; no shipped-game example of a bytecode interpreter in a shader found. + +## 4. Register counts for the trophic graph + +- [src] Aetheria `4354ef2c` (`docs/aetheria-fauna-map.md` B28-B36; `docs/research/2026-10-14-trophic-dag-probe/README.md`) gives layers and costs but no node or register count. B36: "each ring-stencil level multiplies the cost by 9" (krill .26 ms, minnows 3.19 ms = 9 krill, sharks 29.3 ms = 81 krill, Python single evaluation); stencils at most two levels deep (target invariant `trophic-dag-pointwise`). Probe layers: L1-glim, L2-krill (Skein x Holling of glim), L2-minnows (Holling of krill over the point + 8 taps on a 48 u ring), L3-sharks (minnows over a 160 u ring x band), L3-lanterns (slack x krill over a 120 u ring). No tape exists yet. +- [src] `F:\Projects\CultLib\docs\cultmath-interval-ground-cut.md` Pass 7, T7 (`:1677`): in a graph a stencil is "input substitution" (the prey subgraph instantiated per tap with the position input replaced); the builder hash-conses identical nodes, so the 9 taps do not share nodes (their inputs differ). T7(b): site-span ops read an array of wells (per-batch buffer input). T7(c): the lattice weight is evaluated at the site, not the query point. Multi-result ops (cellular) are read through `Pick(node,k)`. +- [src] CultMath already carries a multiplicatively weighted lattice fragment: `packages/cultmath/src/CultMath/Voronoi.cs:151-154` (`weight = o.x*0.5f+0.5f`, `d = max(d2.x,d2.y)*weight`) and an extern `cultmath_apollonian_voronoi_tones` (`:219`). +- [abs] Synthetic probe `probe2.py` (Python; greedy MPR-style slot reuse, free at last use): a 9-tap mean-of-mean tree whose leaf is a 12-op chain with 2 live values. Slots depend on the schedule order, not on op count: + +| stencil levels | ops | slots, depth-first with running accumulator | slots, level-order | +|---|---|---|---| +| 0 | 13 | 2 | 2 | +| 1 | 125 | 3 | 18 | +| 2 | 1,133 | 4 | 149 | +| 3 | 10,205 | 5 | 1,230 | + + A variant that computed all 9 child results before summing gave 10 / 18 / 26 slots at levels 1 / 2 / 3 (+8 per level; first probe script `probe.py`). The tree is a toy (no shared subexpressions, no span inputs, no wide intermediate such as cellular's multi-result values); real counts need the real graph. For scale, MPR's measured counts are 22 to 84 slots for 328 to 6056 clauses (ratio 12.6 to 72.1). +- [abs] Fidget's frequency renumbering and the 255 cap mean the host knows `reg_count` before dispatch and can choose among pipelines (1a). + +## 5. Scope addition: gravity bodies as Voronoi sites (operator sketch) + +### 5a. Diagram variants over a small set of explicit weighted sites +- [web] Wikipedia "Weighted Voronoi diagram" (`https://en.wikipedia.org/wiki/Weighted_Voronoi_diagram`, fetched): multiplicatively weighted = distance divided by a positive weight; boundaries circular arcs and straight segments in the plane; "A Voronoi cell may be non-convex, disconnected and may have holes". Additively weighted = weights subtracted from distance; boundaries hyperbola arcs and straight segments. Power diagram = weights subtracted from squared Euclidean distance. (The page does not state the Apollonius-circle name, convexity or the Delaunay relation; I did not source those.) +- [abs] Equal-potential boundary: for point masses with potential `-G*M/r`, the locus where two bodies' potentials are equal has `r1/r2 = M1/M2`, a multiplicatively weighted Voronoi edge (weight = mass), an Apollonius circle in 2D. The probe's `slack` (weight = gravity radius, citing TD:20 Apollonius circles) is the same family. A sum of potentials `sum G*M_i/r_i` is not a Voronoi cell; its level sets and the points where pulls cancel are different objects (`slack` in the probe is "to first order" `(F2-F1)/|grad(F2-F1)|`). Equal-force (`M/r^2`) boundary: `r1/r2 = sqrt(M1/M2)`, multiplicative with weight `sqrt(M)`. +- [src] Aetheria probe measurements (`docs/aetheria-fauna-map.md` B33-B34): the hard Gabriel graph over six bodies changes its edge set 29 times in 1,400 s (one gas-giant orbit), each a strand appearing or vanishing whole; the soft gate moves at most .00075 per 1/60 s. Largest presence step under the krill carriage: band .0035, slack .013, skein .054, lattice .124, skein-web .493. +- [abs] Pop properties when a site moves: the distance field `F1` and the border measure `F2 - F1` are continuous in site position and weight; what jumps is the discrete label (which site is nearest) and adjacency (which cells touch). For multiplicatively weighted cells a cell can appear, vanish, split or become disconnected as weights change (Wikipedia above). With six bodies, a loop over all sites per point is exact; a lattice-based diagram needs its search window sized to the weight range (Aetheria B34: 5x5 sufficed for a 1.6x boost). +- [web] Cosmic-web analogues: Hidding, van de Weygaert, Vegter, Jones, "Adhesion and the Geometry of the Cosmic Web" (arXiv 1211.5385) and van de Weygaert, "Tessellating the Universe: the Zel'dovich and Adhesion tiling of space" (arXiv 0912.4520, abstract read): the adhesion model (Zel'dovich approximation plus viscosity, Burgers' equation) yields in the zero-viscosity limit a "quasi-Voronoi" tessellation, with the Eulerian/Lagrangian relation tied to Voronoi/Delaunay duality. In that tessellation, 1-dimensional edges correspond to filaments (second most massive elements), vertices to clusters/nodes (most mass), faces/tiles to walls and voids. Voronoi foam kinematic model: matter condenses onto Voronoi walls, edges and nodes (search summary, arXiv 1611.01221 area). The abstract of 0912.4520 does not mention weighted or power diagrams or the distance metric; I found no source tying a weighted Voronoi to the cosmic web. Aetheria already cites a Zel'dovich warp (TD:39). +- [abs] In 3D the Voronoi border set is faces (walls) meeting in edges (filament-like) at vertices; in 2D the borders are edges and the vertices are the node-like points. A rendered "filament" from a diagram's border in 3D is wall-like away from the edges. +- [abs] Anisotropic/flow-stretched metric `d = |A(x-y)|` with constant per-site A elongates cells along A's long axis; a metric varying along the path is a geodesic with no closed form. Not sourced. + +### 5b. Metric as a node parameter +- [doc] Blender Voronoi Texture node (`https://docs.blender.org/manual/en/4.2/render/shader_nodes/textures/voronoi.html`, search summary): Distance metric Euclidean, Manhattan, Chebychev, Minkowski with an Exponent input; exponent 1 = Manhattan, 2 = Euclidean, infinite = Chebychev; manual examples use 0.5, 1, 2, 32 and treat 32 as an approximation of Chebychev; exponent below 1 gives star shapes. Features: F1, F2, Smooth F1, Distance to Edge, N-sphere Radius. +- [src] Cycles implementation (`intern/cycles/kernel/svm/voronoi.h`, fetched from `raw.githubusercontent.com/blender/blender/main/`): `voronoi_distance(a,b,params)` is a template selecting with a chain of `if` on `params.metric`: Euclidean `distance`, Manhattan `reduce_add(fabs(a-b))`, Chebyshev `reduce_max(fabs(a-b))`, Minkowski `powf(reduce_add(power(fabs(a-b), exponent)), 1/exponent)`; unknown metric returns 0. `voronoi_distance_bound` is a cheaper variant that skips the square root (Euclidean) and the final root (Minkowski) for ordering inside the search loops; the true distance is computed once for output. Search windows: 2D 3x3 for F1/F2/edge/n-sphere, 5x5 for Smooth F1; 3D 3x3x3 (27 cells) and 5x5x5 (125) for smooth; 4D 81 and 625. A header comment records the non-smooth windows were reduced from -2..2 to -1..1 for performance. +- [web] Houdini `mx_voronoi` VEX: integer `metric` argument, 0 Euclidean, 1 squared distance, 2 Manhattan (a mirror lists 3 Chebyshev) (`https://www.sidefx.com/docs/houdini/vex/functions/mx_voronoi.html`). Substance 3D Designer "Distance" node: Euclidean, Manhattan, Chebyshev (`https://experienceleague.adobe.com/en/docs/substance-3d-designer/using/substance-graphs/nodes-reference-for-substance-graphs/atomic-nodes/distance`); I did not find the Cells node's own metric options. 3Delight: Minkowski 1 = Manhattan, 2 = Euclidean, 4 rounded square, infinity Chebyshev (`https://documentation.3delightcloud.com/x/TgYuMg`). +- [abs] Cost of a metric switch: Cycles' form is a runtime branch on a parameter uniform across a dispatch (all lanes take the same branch). A Minkowski exponent as a float input costs a `pow` per site pair plus one `pow` for the root; the bound variant avoids the root in loops. In a tape interpreter a constant metric can be chosen at lowering (ops in the graph) or dispatched at run time as one node. + +### Gaps (section 5) +- No source measured the per-evaluation cost of a metric switch; only the Cycles code structure is read. +- No source found for weighted/power diagrams in cosmic-web modelling; the equal-potential = multiplicatively-weighted identity is derived here, not cited. +- No numeric probe of cell geometry was run. diff --git a/docs/research/tape-register-model/probe2.py b/docs/research/tape-register-model/probe2.py new file mode 100644 index 00000000..7023050f --- /dev/null +++ b/docs/research/tape-register-model/probe2.py @@ -0,0 +1,52 @@ +# Register need of a stencil tree under two schedules; MPR-style greedy slot reuse (free at last use). +import sys +LEAF=12 # ops in a shared-input leaf subgraph chain (krill-like): leaf = chain with 3 live temps +def build(levels, taps=9): + # node: ('leaf',) or ('mean', [children]) ; each tap child is an independent subgraph (offset input substituted) + nodes=[] # (operands) + def leaf(): + # leaf subgraph: LEAF ops, each uses previous two results (width 2 live) -> 2 live values + base=len(nodes); nodes.append([]) # input + prev=[base] + for i in range(LEAF): + ops=prev[-2:] + nodes.append(ops); prev.append(len(nodes)-1) + return prev[-1] + def mean(l): + if l==0: return leaf() + acc=mean(l-1) + for _ in range(taps-1): + c=mean(l-1) + nodes.append([acc,c]); acc=len(nodes)-1 + return acc + root=mean(levels) + return nodes,root +def regs(nodes,order): + pos={n:i for i,n in enumerate(order)} + last={} + for n in order: + for o in nodes[n]: last[o]=max(last.get(o,-1),pos[n]) + live=0;peak=0;free=0;alloc=0 + for n in order: + # free operands whose last use is here + for o in set(nodes[n]): + if last[o]==pos[n]: free+=1 + if free: free-=1 + else: alloc+=1 + if n not in last: free+=0 + return alloc +def dfs_order(nodes): + return list(range(len(nodes))) # construction order is postorder DFS +def bfs_order(nodes,root): + # level order by depth from root (reverse), a bad schedule + depth={root:0} + users={} + for n,ops in enumerate(nodes): + for o in ops: users.setdefault(o,[]).append(n) + d=[0]*len(nodes) + for n in range(len(nodes)-1,-1,-1): + d[n]=max([d[u]+1 for u in users.get(n,[])],default=0) + return sorted(range(len(nodes)),key=lambda n:(-d[n],n)) +for L in range(0,4): + nodes,root=build(L) + print(L,"ops",len(nodes),"slots_dfs",regs(nodes,dfs_order(nodes)),"slots_levelorder",regs(nodes,bfs_order(nodes,root)))