docs(geometry): record virtual draw evidence

Author: proggeramlug

docs/evidence/issue-131-virtual-draw-emission-v1.json

@@ -0,0 +1,109 @@
+{
+  "schema": "bloom-virtual-draw-emission-evidence-v1",
+  "issue": 131,
+  "qualified_revision": "92daf0eae684ca5a0b359cc8c991f728ffb83600",
+  "gpu_abi": {
+    "draw_command_bytes": 16,
+    "emission_state_bytes": 48,
+    "dispatch_command_bytes": 12,
+    "float32_vertex_bytes": 72,
+    "quantized_vertex_bytes": 32,
+    "local_index_bytes": 1,
+    "draw_capacity_matches_selected_capacity": true,
+    "first_instance_addresses_selected_record": true
+  },
+  "construction_limits": {
+    "storage_buffer_count_checked": true,
+    "command_buffer_size_checked": true,
+    "storage_binding_size_checked": true,
+    "workgroup_size_checked": true,
+    "dispatch_dimension_checked": true,
+    "selector_identity_checked": true
+  },
+  "whole_batch_fallback": {
+    "selected_attempts": 4,
+    "selected_capacity": 2,
+    "selected_overflow": 2,
+    "published_draw_count": 0,
+    "published_dispatch_x": 0,
+    "batch_fallback": 1,
+    "emitted_draws": 0,
+    "emitted_triangles": 0,
+    "invalid_missing_counter_saturates": true,
+    "capacity_mismatch_fails_closed": true
+  },
+  "resident_ancestor_request_overflow": {
+    "request_attempts": 2,
+    "request_capacity": 1,
+    "request_overflow": 1,
+    "resident_ancestor_draws": 2,
+    "batch_fallback": 0
+  },
+  "metal_decode_oracle": {
+    "backend": "metal",
+    "selected_clusters": 4,
+    "decoded_corners_per_format": 12,
+    "encodings": ["float32", "quantized32"],
+    "position_match": true,
+    "normal_match": true,
+    "tangent_match": true,
+    "uv0_match": true,
+    "uv1_match": true,
+    "color_match": true,
+    "local_index_match": true,
+    "maximum_test_tolerance": 0.00001
+  },
+  "metal_indirect_raster_oracle": {
+    "commands": 4,
+    "triangles": 4,
+    "target_format": "Rgba8Uint",
+    "readback_red": [1, 2, 3, 4],
+    "first_instance_exact": true,
+    "same_encoder_after_compute": true
+  },
+  "default_path": {
+    "engine_state_fields_added": 0,
+    "renderer_fields_added": 0,
+    "model_manager_fields_added": 0,
+    "passes_added": 0,
+    "draws_added": 0,
+    "buffers_constructed": 0,
+    "bindings_added": 0,
+    "shader_branches_added": 0,
+    "pixels_changed": false
+  },
+  "qualification": {
+    "quick_lane": "pass",
+    "quick_lane_seconds": 51,
+    "shared_unit_tests_passed": 401,
+    "shared_unit_tests_ignored": 1,
+    "virtual_geometry_tests_passed": 22,
+    "golden_tests_passed": 59,
+    "golden_tests_ignored": 2,
+    "render_target_tests_passed": 4,
+    "visibility_parity_tests_passed": 2,
+    "visibility_runtime_tests_passed": 2,
+    "cooker_tests_passed": 29,
+    "quality_governance_tests_passed": 39,
+    "visual_fault_engine_tests_passed": 3,
+    "strict_production_clippy": "pass",
+    "format": "pass",
+    "contracts_ffi_file_ratchet": "pass",
+    "wasm_no_default_check": "pass",
+    "wasm_models3d_check": "pass",
+    "native_models3d_check": "pass",
+    "no_default_optional_dependency_boundary": "pass"
+  },
+  "acceptance_boxes_changed": 0,
+  "remaining": [
+    "current and previous transform plus material remap ABI",
+    "visibility and exact PBR material composition",
+    "bounded submission on adapters without indirect-count support",
+    "conservative previous-frame Hi-Z occlusion",
+    "asynchronous request feedback and artifact streaming",
+    "crack and temporal image corpus",
+    "per-MRT parity",
+    "10-million-triangle stress and total GPU timing",
+    "integrated and discrete adapter qualification"
+  ]
+}

docs/evidence/issue-131-virtual-draw-emission-v1.md

@@ -0,0 +1,117 @@
+# Issue #131 virtual draw emission and raw-page decode v1 evidence
+
+This checkpoint qualifies Bloom's first bounded virtual-cluster indirect
+stream and cooked-page GPU decoder at revision
+`92daf0eae684ca5a0b359cc8c991f728ffb83600`. It proves that selected resident
+clusters can become executable indirect work without expanding their cooked
+vertices. It does not yet register a production render path or change ordinary
+scene pixels.
+
+## Fixed ownership and ABI
+
+`GpuVirtualDrawEmitter` is constructed explicitly from one hierarchy selector.
+It rejects another selector, validates command-buffer and workgroup limits
+before allocation, and owns three fixed resources:
+
+| Record | Exact bytes | Purpose |
+|---|---:|---|
+| non-indexed draw command | 16 | triangle vertex count and selected-record address |
+| emission state/count | 48 | indirect count, whole-batch fallback, and counters |
+| compute dispatch command | 12 | bounded indirect emission dispatch |
+
+The command capacity exactly equals the selector's selected-cluster capacity.
+The prepare pass reads selection counters and publishes a bounded draw count and
+dispatch. The emit pass writes one command per admitted cluster; its
+`first_instance` is the matching selected-record index. Triangle counts come
+only from the strictly validated archive and remain bounded by the 4 MiB page
+ceiling.
+
+The state count lives at byte zero for `multi_draw_indirect_count`. That optional
+feature is not assumed for production activation: adapters without it must use
+a separately qualified bounded submission strategy or remain on compatibility
+rendering. This checkpoint's Metal raster oracle uses the supported fixed-count
+form with the fixture's known four commands.
+
+## Whole-batch safety
+
+The prepare shader suppresses the complete virtual batch when selection
+overflowed, the selector observed an invalid record, a current selected page is
+missing, or selected/count capacities disagree. Invalid-plus-missing telemetry
+saturates instead of wrapping. No partial geometry is published as a valid
+batch.
+
+The real Metal overflow fixture attempts four leaf selections against a
+two-record capacity and proves:
+
+- `selected_count = 4`, `selected_overflow = 2`;
+- `draw_count = 0`, `batch_fallback = 1`;
+- indirect dispatch x is zero;
+- emitted draw and triangle counters remain zero.
+
+Request overflow has different semantics. With only complete resident middle
+groups available, two leaf-page requests are attempted against one request
+slot, but the two resident ancestors remain a complete valid fallback batch.
+The emitter therefore publishes two draws with no batch fallback. This prevents
+feedback pressure from creating a visible hole.
+
+## Raw cooked-page vertex pulling
+
+The common WGSL decoder consumes the fixed physical page buffer directly. It
+supports both validated archive encodings:
+
+- version 1: 72-byte Float32 position, normal, tangent, UV0, UV1, and color;
+- version 2: 32-byte quantized AABB-local position, octahedral normal/tangent,
+  binary16 UV0/UV1, UNORM8 color, handedness, and the missing-tangent flag.
+
+It reads each cluster's local `u8` corner index, page-local vertex/index offsets,
+stride, AABB, and mesh encoding from the production GPU tables. The Metal probe
+selects four resident leaves, decodes all twelve indexed corners for each
+format, and compares every output lane plus selected/cluster/corner/local-index
+identity. Float32 and quantized positions, normals, tangents, UV sets, and
+colors match their CPU-authored values within `1e-5`.
+
+## Executable Metal indirect proof
+
+A real render pass consumes the GPU-written command buffer in the same command
+encoder after hierarchy selection and emission. Four indirect triangles write
+four distinct `Rgba8Uint` pixels. Readback is exactly `[1, 2, 3, 4]` in the red
+channel, proving that every command executed and that `first_instance` retained
+the expected selected-record index.
+
+Together with the independent traversal oracle, the focused virtual-geometry
+suite now has 22 passing tests covering strict load, fixed residency, retirement,
+GPU table identity, LOD/frustum/cone selection, missing-page fallback, bounded
+overflows, camera cuts and motion, both vertex formats, command emission, and
+actual indirect rasterization.
+
+## Regression boundary and gates
+
+The pool, selector, emitter, and decoder have no `Renderer`, `EngineState`,
+`ModelManager`, frame-graph, scene, FFI, or ordinary glTF owner. The shipping
+path therefore constructs none of their buffers, pipelines, bindings, passes,
+draws, or branches. Existing rendering pixels, allocations, and steady-state
+work are unchanged.
+
+The complete governed quick lane passed in 51 seconds on the qualified tree:
+
+- 401 shared unit tests passed and one existing hot-reload test was ignored;
+- the negotiated real-device test passed;
+- 59 golden tests passed and two hardware-specific tests were ignored;
+- four render-target tests, two visibility parity tests, and both visibility
+  runtime tests passed;
+- strict correctness/suspicious/performance Clippy, formatting, FFI/schema
+  parity, CI contracts, and file-size governance passed;
+- 39 quality-governance tests, three visual fault-engine tests, and 29 cooker
+  tests passed;
+- baseline wasm, native `models3d`, and wasm `web,models3d` checks passed;
+- the no-default dependency graph still contains neither
+  `bloom-geometry-format` nor SHA-256.
+
+## Remaining acceptance work
+
+No #131 acceptance checkbox changes at this checkpoint. Selected virtual work
+still needs current/previous transform and material-remap records, integration
+with #27's exact PBR visibility composition, a bounded non-indirect-count path,
+conservative previous-frame Hi-Z, asynchronous request readback/streaming,
+crack and temporal image corpora, 10-million-triangle stress, per-MRT parity,
+total GPU timings, and integrated/discrete adapter qualification.

docs/virtualized-geometry.md

@@ -1,11 +1,11 @@
 # Virtualized geometry
 
-Bloom's virtualized-geometry work is opt-in and staged. The offline #131
-milestones establish a deterministic meshlet/page contract; the first runtime
-milestone adds one shared strict reader, immutable archive ownership, and a
-fixed-budget residency plan with deterministic ancestor fallback. It does not
-yet enable a new renderer, change ordinary glTF loading, or claim
-Nanite-equivalent GPU traversal/streaming.
+Bloom's virtualized-geometry work is opt-in and staged. The #131 milestones now
+cover the deterministic meshlet/page contract, strict runtime loading,
+fixed-budget GPU residency, projected-error GPU hierarchy selection, raw-page
+vertex decoding, and bounded indirect draw emission. They still do not enable
+a production renderer, change ordinary glTF loading, or claim complete
+Nanite-equivalent streaming, occlusion, material, or visibility integration.
 
 ## Cook and inspect
 
@@ -234,27 +234,54 @@ resident state, and pinned state. Telemetry distinguishes allocated GPU bytes,
 resident slot bytes, useful payload bytes, pinned/retiring slots, frame and
 lifetime upload/eviction counts, denials, and exact/ancestor fallback results.
 
+`GpuVirtualHierarchySelector` consumes fixed 128-byte instance records and
+walks each atomic hierarchy group on the GPU. Projected pixel error controls
+LOD. Frustum and transform-safe normal-cone tests reject invisible work;
+non-uniform or sheared transforms conservatively disable only cone rejection.
+Refinement requires the complete child group to be resident. Otherwise the
+selector keeps the resident ancestor and writes bounded page requests. Camera
+cuts and instance motion consume no prior selection state.
+
+`GpuVirtualDrawEmitter` converts the bounded selected table into compact
+16-byte non-indexed indirect commands. `first_instance` addresses the matching
+selected-cluster record for vertex pulling. Selection overflow, an invalid
+record, or a missing current page publishes a zero draw count for the whole
+virtual batch; request overflow remains safe because the complete resident
+ancestor batch is still selected. Construction validates command-buffer and
+compute-dispatch limits up front.
+
+The shared WGSL page decoder reads local `u8` triangle indices directly from
+the fixed physical pool and reconstructs every version 1 Float32 or version 2
+quantized vertex lane. A Metal compute oracle reads the real mesh, cluster,
+selection, and physical-page buffers. A separate Metal render oracle consumes
+the emitted indirect commands and proves the exact `first_instance` values.
+An eventual production consumer must use indirect-count support or provide a
+separately qualified bounded fallback; unsupported adapters remain on the
+compatibility renderer.
+
 The renderer integration remains deliberately explicit:
 
 - zero production render passes, draws, buffers, bindings, allocations, or
   shader branches unless a caller explicitly constructs the GPU pool;
 - zero changes to existing immediate-mode or glTF selection and pixels;
 - no silent `.bgeo` replacement of an ordinary model;
-- no asynchronous file/store IO, feedback readback, or traversal dispatch yet.
+- no asynchronous file/store IO, request-feedback scheduling, production
+  visibility raster, or virtual-geometry material composition yet.
 
-The next #131 runtime milestones are asynchronous #136 index resolution and
-request feedback, projected-error hierarchy traversal, and integration with
-#27's visibility/material path and #28's GPU-driven submission. The
-compatibility renderer remains responsible for unsupported and
-not-yet-virtualized content throughout that work.
+The next #131 runtime milestones are temporal instance/material records,
+integration with #27's visibility/material path, a non-indirect-count fallback,
+conservative previous-frame Hi-Z, and asynchronous #136 index resolution plus
+request feedback. The compatibility renderer remains responsible for
+unsupported and not-yet-qualified content throughout that work.
 
 The default version 1 artifact remains byte-identical to the qualified
 leaf-only milestone (`parent` and `first_child` absent, both relation counts
 zero, level/error zero). Opt-in hierarchy artifacts populate those formerly
 reserved fields without changing the 128-byte cluster record. Opt-in packed
 vertices use format version 2 so a version 1 reader can never reinterpret the
-32-byte stride as float32. Runtime streaming remains disabled until residency,
-traversal, occlusion, and fallback milestones are independently qualified.
+32-byte stride as float32. Runtime activation remains disabled until temporal
+data, visibility/material composition, streaming feedback, occlusion, and
+platform fallback milestones are independently qualified.
 
 ## Qualification
 
@@ -303,3 +330,7 @@ loose-store qualification are recorded in
 The fixed physical GPU pool, stable ID/page-table ABI, bounded frame work, and
 Metal readback proof are recorded in
 `docs/evidence/issue-131-gpu-page-pool-v1.{md,json}`.
+GPU hierarchy selection and independent CPU parity are recorded in
+`docs/evidence/issue-131-gpu-hierarchy-traversal-v1.{md,json}`. Raw-page decode,
+bounded indirect emission, and executable Metal raster proof are recorded in
+`docs/evidence/issue-131-virtual-draw-emission-v1.{md,json}`.