From 295e9be5736a92425b4281047576d959bbc1d5b2 Mon Sep 17 00:00:00 2001 From: Jeremy Date: Wed, 21 Oct 2026 00:25:44 +0000 Subject: [PATCH 1/3] docs(research): expand GPU grid algorithms and independent backends --- .../gpu-grid-research/execution-ledger.md | 54 ++ docs/plans/gpu-grid-research/issues.md | 13 + docs/plans/gpu-grid-research/plan.md | 94 ++++ .../reports/draft-pr-handoff.md | 49 ++ .../reports/independent-review.md | 147 +++++ .../reports/standards-and-scope.md | 34 ++ .../chapters/37-gpu-grid-algorithms.md | 356 ++++++++++++ .../38-independent-cuda-and-vulkan.md | 520 ++++++++++++++++++ .../chapters/39-gpu-validation-and-roadmap.md | 355 ++++++++++++ .../gpu-cpu-oracle-inventory.md | 428 ++++++++++++++ .../implementation/gpu-research-plan.md | 116 ++++ .../sources/gpu-algorithm-bibliography.md | 87 +++ .../sources/gpu-backend-bibliography.md | 270 +++++++++ 13 files changed, 2523 insertions(+) create mode 100644 docs/plans/gpu-grid-research/execution-ledger.md create mode 100644 docs/plans/gpu-grid-research/issues.md create mode 100644 docs/plans/gpu-grid-research/plan.md create mode 100644 docs/plans/gpu-grid-research/reports/draft-pr-handoff.md create mode 100644 docs/plans/gpu-grid-research/reports/independent-review.md create mode 100644 docs/plans/gpu-grid-research/reports/standards-and-scope.md create mode 100644 docs/research-book/chapters/37-gpu-grid-algorithms.md create mode 100644 docs/research-book/chapters/38-independent-cuda-and-vulkan.md create mode 100644 docs/research-book/chapters/39-gpu-validation-and-roadmap.md create mode 100644 docs/research-book/implementation/gpu-cpu-oracle-inventory.md create mode 100644 docs/research-book/implementation/gpu-research-plan.md create mode 100644 docs/research-book/sources/gpu-algorithm-bibliography.md create mode 100644 docs/research-book/sources/gpu-backend-bibliography.md diff --git a/docs/plans/gpu-grid-research/execution-ledger.md b/docs/plans/gpu-grid-research/execution-ledger.md new file mode 100644 index 00000000..cfdd53cb --- /dev/null +++ b/docs/plans/gpu-grid-research/execution-ledger.md @@ -0,0 +1,54 @@ +# GPU research execution ledger + +## 2026-10-21 — source and isolation milestone + +Preserved clean original `work` at `6ba41c7`. Fetched `origin/main` to +`fd7c0f0aee63bc8710df99f9bb01154e7619b974` and created isolated research branch +and worktree recorded in the [plan](plan.md). No product/FEM/shared book files +were edited. Tracked and ancestor `AGENTS.md` and repository `.agents/skills` +were absent; workspace `.agents` and `.codex` were empty. + +Read current source/output policy, book README/builders/education workflow, +existing scheduling/memory/validation material, actual solver source and current +research audit. Read fetched Coding-Standards authority at `0fc92a2` as detailed +in the standards report. Launched bounded GPT-6.1-sol high-reasoning CPU inventory, +GPU algorithm/source and independent backend/source investigations, each owning +new disjoint paths. No scientific/hardware workload or package installation ran. + +## 2026-10-21 — draft, review and handoff milestone + +Completed seven new GPU book files: CPU inventory, reading/source map, three +substantial chapters and two bibliographies, about 22,000 words together. Parent +confirmed FEM owns Chapters 30–36 at `9489aa3`; only the new GPU chapters moved +to 37–39 and their owned crosslinks changed. Existing book/shared TOC/README, +FEM/product source and Appendix C are untouched. Fetched main again at completion; +it remained `fd7c0f0`. The original `work` checkout remained clean. + +Independent GPT-6.1-sol high-reasoning review read CPU source and reopened +deciding papers/vendor specifications/pinned upstream code/licenses. It found +pressure acceptance ordering, residual-replacement flag initialization, +final-budget PCG failure classification, gauge comparison and a draft-link issue. +All were repaired and reread with no unresolved deciding findings. An additional +CPU-source review corrected numerical commit ownership, observed-cancel races, +recording-profile admission and the description of historical cross-wall +limitations. See the [independent review](reports/independent-review.md) for +its exact evidence and retrieval limits. Review is documentation/source evidence, +not numerical, hardware or performance qualification. + +Final nonnumerical staged checks covered 13 new Markdown files and 174 local links, with no +missing targets, stale GPU chapter paths, title/newline/trailing-whitespace, +fence or math-delimiter issues. Scoped Pandoc conversion of only the three +chapters to temporary HTML succeeded: 11 tables, 25 math spans and 9 code blocks. +This is a parse check, not equation rendering, Word/table layout or browser +qualification. New chapter code fences follow the manuscript builder's `~~~` +convention. Staged scope/whitespace are checked again at the local commit boundary. + +No solver/tests/scientific/GPU/hardware workload, full book build, render, +installation or security change ran. Source fetch limitations are recorded in +the bibliographies/review; the supported narrow claims have primary evidence. +No auto-review rejection or uncertain mutation occurred. No push/PR/merge, +deployment or publication occurred. The [parent handoff](reports/draft-pr-handoff.md) +contains concrete draft PR text and the workflow-scope decision: normal push/PR +can trigger scientific jobs, so parent must coordinate the later action without +bypassing gates. The isolated committed branch/worktree remain retained under +the explicit review/retirement contract in the plan. diff --git a/docs/plans/gpu-grid-research/issues.md b/docs/plans/gpu-grid-research/issues.md new file mode 100644 index 00000000..9aca722a --- /dev/null +++ b/docs/plans/gpu-grid-research/issues.md @@ -0,0 +1,13 @@ +# GPU research issues and dispositions + +| ID | Finding | Disposition / owner | +| --- | --- | --- | +| GPU-D01 | Parent confirms FEM occupies Chapters 30–36 at `9489aa3` | Resolved: only new GPU files moved to 37–39; parent owns shared integration, existing material untouched | +| GPU-D02 | Full book/Pages build includes scientific/reference/native qualification; Word builder lacks dedicated Markdown-table rendering | No build run under this task. Parent must inspect/authorize full scope and perform output review separately | +| GPU-D03 | No installed device/API/compiler evidence has been collected | Explicitly outside scope; future implementations must query and qualify each backend/device configuration | +| GPU-D04 | General evolving 3D free-surface, arbitrary mesh stress and coupled physics are not supplied by the bounded CPU carrier | Numerical research/qualification precedes port; Chapters 37–39 retain model-family gates | +| GPU-D05 | Moving official documentation is not an immutable SDK/toolchain snapshot | Bibliographies record consulted version/access date; future implementation pins actual environment | +| GPU-D06 | Existing host-slice pressure provider cannot establish whole-step device residency | Separate pressure-only diagnostic prototype from a later resident backend; future implementation owner | + +No blocked or denied operation has occurred at this checkpoint. Publication, +hardware and scientific qualification are excluded work, not uncertain outcomes. diff --git a/docs/plans/gpu-grid-research/plan.md b/docs/plans/gpu-grid-research/plan.md new file mode 100644 index 00000000..9c728bc5 --- /dev/null +++ b/docs/plans/gpu-grid-research/plan.md @@ -0,0 +1,94 @@ +# GPU grid-fluid research book expansion + +Status: Accepted documentation deliverable. Production implementation is excluded. +Current phase: terminal local research handoff; parent integration remains separate. +Next slice: none for this worker's accepted documentation scope. + +## Objective, constraints and admission + +Expand Puma's existing research book with substantial GPU algorithms grounded in +Rheon's current CPU solver, and compare independently usable CUDA and Vulkan +compute paths. Deliver source inventory, primary-source map, chapters, comparative +architecture and reviewable future validation/implementation stages. User +authorization admits this research work, isolation and independent review; it +does not admit production backend work, merge, deployment or auto-publication. + +Fetched main/base: `fd7c0f0aee63bc8710df99f9bb01154e7619b974`. +Branch: `research/gpu-grid-algorithms-20261021`. +Worktree: `/workspace/rheon-gpu-research`. +Owner: GPU research worker; parent owns integration and later draft-PR +coordination. Original `/workspace/Rheon` checkout remains preserved. PR145, +PR146 and Lagrangian FEM are outside scope. No shared numbering, TOC, build or +existing reference edits are authorized by this worker's ownership. + +Keep source/field identities, original credit, historical evidence and failed +criteria. Do not run scientific/GPU/hardware workloads or install toolchains, +drivers or modify security settings. Only nonnumerical documentation, link, +format and source-scope checks may execute. Current primary sources must be read; +unavailable source or capabilities stay explicit. No benchmark or proof is +fabricated. GPT-6.1-sol at high reasoning is used for delegated CPU/source and +technical review work as requested. + +## Owned artifacts and decisions + +The [book source map](../../research-book/implementation/gpu-research-plan.md) +lists the exact new book paths. This plan directory owns plan, execution ledger, +issues and review/standards reports. Only these new GPU paths may change. + +The CPU inventory owns claims about existing source. Chapter 37 owns algorithm +comparison, Chapter 38 owns independent backend mapping, Chapter 39 owns future +validation and roadmap. Bibliographies own source provenance. No implementation +contract is silently admitted by these prospective documents. + +Composed-design review: applicable to the proposed multi-backend architecture; +documentation changes no runtime composition. Chapter 39 and the following probe +record the candidate design to be challenged in later implementation planning. + +| Architecture probe | Research disposition | +| --- | --- | +| Independent what/who/how/when/where/why | Numerical owners select equations/field stages; CUDA and Vulkan independently schedule on their devices; sessions commit complete candidates; recorders/renderers consume accepted versions. | +| Necessary and accidental interleavings | Numerical stage dependencies and publication fences are necessary. Renderer timing, archive serialization, API objects and optional interop must not define the physics. | +| Caller/composition knowledge | Select admitted problem, backend capability/arithmetic policy, resource/cancel/output budget. Do not expose launch/block/barrier details through physical controls. | +| Representative change propagation | New preconditioner affects method/qualification; archive codec affects recording; Vulkan barrier fix affects Vulkan execution; new physical traction affects numerical contracts/all implementing backends/oracles. | +| Stable values versus hidden knowledge | Grid association, units, interval/generation, geometry and acceptance are stable meanings. API handles, descriptor/graph lifetimes and native layouts remain backend-owned. | +| Replaceability and deletion | Each backend can be disabled/removed independently; CPU remains reference. Vulkan must build/run without CUDA; CUDA compute must not require Vulkan. No shared opaque GPU execution object is proposed. | +| Cumulative cost and alternatives | Start with dense Jacobi PCG; defer sparse metadata, graph optimization, multigrid and optional interop until correctness/measurement. Compare pressure-only host provider prototype with resident full step and disclose transfer costs. | +| Simplest coherent artifact | Three complementary research chapters plus source inventory/bibliographies. No generic runtime registry, new archive, shader/kernel code or production abstraction is introduced. | + +## Milestones and acceptance + +| Milestone / write set | Goal and gate | State | +| --- | --- | --- | +| M0 / new source inventory and plan | Fetch actual main; inspect guidance, book gaps and build scope; preserve unrelated work | Complete | +| M1 / new chapters and bibliographies | Equations/assumptions, pseudocode, comparative independent backends, capabilities, precision, lifecycle and roadmap, grounded in primary sources | Complete | +| M2 / new review report and owned repairs | Independent technical and source review; address deciding findings; nonnumerical local links/format/scope checks | Complete; [review](reports/independent-review.md) has no unresolved deciding findings | +| M3 / ledger and parent handoff | Exact artifacts/branch/commit and evidence limits; coordinate numbering/shared integration and a later draft PR through parent | Handoff prepared; [draft PR text and constraints](reports/draft-pr-handoff.md), no remote publication | + +Objective acceptance is documentation-level only: (A) actual CPU source and +capability limits inventoried, (B) substantial principles-first GPU algorithms +and independent mappings, (C) current primary sources with provenance and license +boundaries, (D) explicit validation/benchmark/lifecycle roadmap, and (E) +independent review and nonnumerical checks. Current acceptance: satisfied locally. +GPU speed, hardware support, numerical parity/convergence and book render/build +are prospective, not failed or passed objective claims. + +Re-plan when source main changes relevant CPU semantics, another worker claims +numbering/files, source evidence contradicts an algorithm/device claim, or the +scope asks for executable backend/qualification. Stop an affected denied action +and report exact target/reason/outcome. Continue independent authorized work. + +## Retained review resource and terminal disposition + +The branch/worktree are retained for parent source review and draft-PR +coordination, not integrated. Parent owns final numbering/shared edits and +eventual integration or rejection. No automatic rebase or history rewrite; +refresh relevant main differences read-only and reconcile through a separately +reviewed change. Supported consumer: parent/reviewer only. Retirement condition: +parent records integration/rejection and disposes of the clean isolated resource +under repository lifecycle rules. No merge, push or PR is implied by local +documentation acceptance. + +See [ledger](execution-ledger.md), [issues](issues.md) and +[standards report](reports/standards-and-scope.md), +[independent review](reports/independent-review.md) and +[parent draft-PR handoff](reports/draft-pr-handoff.md). diff --git a/docs/plans/gpu-grid-research/reports/draft-pr-handoff.md b/docs/plans/gpu-grid-research/reports/draft-pr-handoff.md new file mode 100644 index 00000000..b169ab34 --- /dev/null +++ b/docs/plans/gpu-grid-research/reports/draft-pr-handoff.md @@ -0,0 +1,49 @@ +# Parent handoff for a later draft PR + +No remote push or PR creation was performed by this worker. Parent coordination +and review own the next publication decision. A push/PR can trigger the existing +research-book workflow's scientific/native qualification; creating a draft PR +does not itself prevent those jobs. The current task prohibits those workloads, +so parent must resolve that scope before using the normal remote workflow. This +note does not propose bypassing or disabling any gate. + +Suggested title: `docs(research): expand GPU grid algorithms with independent CUDA and Vulkan` + +Suggested description: + +> Rheon's book has CPU memory/scheduling guidance but lacks a substantial GPU +> treatment grounded in its actual MAC solver. Add Chapters 37–39 covering +> compatible grid algorithms, independently usable CUDA/Vulkan compute paths, +> and source-identified CPU validation, residency, lifecycle and benchmark +> milestones. Include the current CPU inventory and verified primary-source +> bibliographies. Puma's author credit is preserved. +> +> The initial port target remains the bounded sealed constant-density MAC box. +> Particle appearance voxels, reduced surface/obstacle models and future general +> 3D coupled physics are explicitly distinguished. No backend, dependency, product +> source, proof or numerical threshold changes. +> +> Validation: independent technical/source review with deciding findings repaired; +> local source/link/fence/math-delimiter/whitespace checks and isolated Pandoc +> parsing only. No full book build, render, scientific/hardware workload or GPU +> qualification ran. All performance and device-support claims remain prospective. + +## Exact integration boundaries + +Use the [source map](../../../research-book/implementation/gpu-research-plan.md) +for every authored book path. Chapters 30–36 belong to the parallel FEM draft +`9489aa3c1bd6652f9ab1874fce904853b9063dc3` on +`research/lagrangian-fem-book-20261021`. Only the new GPU chapters were renamed +37–39 after this parent instruction. This branch does not include or modify FEM. + +Shared book README, Chapters 11/29, fitted-height/fitted-ALE guides, TOC/build +scripts and Appendix C are untouched. Parent should integrate the two branches' +new material serially and decide shared navigation/coverage/reference placement. +The current builder discovers new chapter files, but reading/PDF/Word and source +bibliography distribution still require their normal scope/freshness/output review. + +Review this research branch against its admitted base +`fd7c0f0aee63bc8710df99f9bb01154e7619b974`, then inspect newer main for relevant +CPU/source changes. Do not infer numerical or lifecycle compatibility after a +changed CPU oracle merely because Markdown applies cleanly. No rebase, merge, +deploy or publication is authorized by this handoff. diff --git a/docs/plans/gpu-grid-research/reports/independent-review.md b/docs/plans/gpu-grid-research/reports/independent-review.md new file mode 100644 index 00000000..c68af395 --- /dev/null +++ b/docs/plans/gpu-grid-research/reports/independent-review.md @@ -0,0 +1,147 @@ +# Independent technical and primary-source review + +Review date: 2026-10-21. CPU baseline: `fd7c0f0aee63bc8710df99f9bb01154e7619b974`. +Research branch: `research/gpu-grid-algorithms-20261021`. +Reviewer: separate technical/source review worker; original book authorship remains Puma. + +**Final disposition: no unresolved deciding findings in the reviewed GPU +documentation after the repairs below were reread.** This is source/documentation +acceptance only; no implemented GPU method or numerical result is qualified. + +This review covers the new CPU inventory, reader/source plan, Chapters 37–39, +and the algorithm/backend bibliographies. The initial drafts were numbered +30–32; parent coordination moved only the new GPU files to 37–39 to preserve +the separately owned FEM chapter range. This report records source inspection +and document checks, not numerical, device, performance or publication +qualification. The reviewer changed only this report; authors/parent made the +document repairs described below. + +## Ranked deciding findings and repairs + +No critical or high-severity source/API finding was identified. The following +medium-severity technical issues and low-severity navigation issue were found +in drafts and sent to their owners. Locations below use the final filenames; +line numbers identify the inspected repaired passage unless explicitly labeled +as original draft lines. + +1. **Medium — preserve pressure acceptance before stored-field acceptance.** + In [Chapter 37](../../../research-book/chapters/37-gpu-grid-algorithms.md), + original draft line 181 made PCG acceptance require “stored correction + checks.” That mixed the pressure candidate contract with the separate full + step correction and could change solver iteration/stopping behavior or + accidentally correct the candidate repeatedly. `src/pressure.rs:563` accepts + the freshly recomputed full residual and its predicted divergence; + `src/simulation.rs:944` then solves, corrects stored f32 velocity, and checks + actual divergence at `src/simulation.rs:959`. Correction: PCG returns + pressure success from true-residual/predicted-divergence gates; the step + performs the one stored correction and independent gate afterward. + **Resolved and reread:** Chapter 37 lines 169, 183–184 and 197 now state that + separation. Chapter 39 lines 70–71 also show the correct sequence. + +2. **Medium — reset the residual-replacement direction flag every iteration.** + The Chapter 37 draft set `restart_direction = true` only on replacement, + then read the flag unconditionally; it never established false on the + ordinary path. An implementation of that sketch could read an undefined + flag or retain true forever after one replacement, degenerating the + intended PCG direction policy. `src/pressure.rs:614` establishes + `replaced = false` inside each iteration before a candidate check. + **Resolved and reread:** Chapter 37 line 173 initializes the flag inside the + loop, before any branch. + +3. **Medium — return iteration exhaustion before preparing a nonexistent next iteration.** + Chapter 37 lines 185–189, as first rechecked, still preconditioned the + residual and validated `gamma_next` after an unsuccessful final-budget true + check. The CPU returns `IterationLimit` at `src/pressure.rs:625` before the + next preconditioner application at line 631. The draft could mask exhaustion + with a later breakdown and perform unnecessary work. Correction: return + iteration exhaustion immediately after the unsuccessful final candidate + check, before preconditioning. **Resolved and reread:** Chapter 37 line 187 + returns `IterationLimit` before the preconditioner at line 188. + +4. **Medium — align two pressure fields with their own gauge values.** + The [Chapter 39](../../../research-book/chapters/39-gpu-validation-and-roadmap.md) + draft instructed subtracting the same pinned-cell value from both fields. + Subtracting a common constant leaves their difference unchanged and does + not remove a relative gauge offset. Subtract each field's own value at the + same pin, or shift the candidate to the reference pin. The sealed CPU + convention is pressure cell zero fixed to zero (`src/pressure.rs:409`, + `src/pressure.rs:555`). **Resolved and reread:** Chapter 39 lines 112–116 now + specify each field's own pinned value and retain the restriction against + arbitrarily shifting pressure-anchored boundaries. + +5. **Low — link the actual independent-backend chapter filename.** + Chapter 37 original draft line 7 linked the nonexistent + `31-cuda-vulkan-backends.md`. **Resolved and reread:** the link now targets + `38-independent-cuda-and-vulkan.md`, including the final renumbering. The + final scan also found no missing local relative-link destinations. + +## Deciding CPU-source checks + +The pressure/operator and publication claims are the deciding boundary for +this research. Direct source reads confirmed the following; no test was run. + +| Claim reviewed | Deciding source evidence | Disposition | +| --- | --- | --- | +| Named-axis staggered f32 velocities/tracer, x-fast layout, six f64 pressure vectors and baseline payload | `src/geometry.rs:81`, `src/geometry.rs:235`, `src/geometry.rs:281`; `src/simulation.rs:158`; `src/pressure.rs:235` | Consistent. Payload is managed array storage, not total process/device memory. | +| Integrated pressure sign/coefficient and physical correction | `src/operator.rs:57`, `src/operator.rs:318`, `src/operator.rs:399`, `src/operator.rs:568` | Consistent with `K=B diag(area/(rho*distance)) B^T`, negative outward-flux RHS and subtraction of physical pressure gradient. | +| Zero guess, sealed compatibility, eliminated gauge, full ungauged true residual and conjunction of residual/predicted divergence | `src/pressure.rs:374`, `src/pressure.rs:545`, `src/pressure.rs:555`, `src/pressure.rs:563`, `src/pressure.rs:586` | Consistent. No mean-shift compatibility repair or pressure warm start is invented. | +| Actual candidate stage order and stored-velocity checks | `src/simulation.rs:655`, `src/simulation.rs:792`, `src/simulation.rs:818`, `src/simulation.rs:860`, `src/simulation.rs:944`, `src/simulation.rs:959`, `src/simulation.rs:984`, `src/simulation.rs:1006`, `src/simulation.rs:1080` | Consistent after acceptance-order repair. Tracer uses corrected candidate velocity; its integral is diagnostic. | +| Ordered compensated pressure reductions and SGS cross-row dependencies | `src/pressure.rs:127`, `src/pressure.rs:142`, `src/pressure.rs:406`; `src/sgs.rs:1`, `src/sgs.rs:20`, `src/sgs.rs:56` | Consistent. Atomic storage contains f64 bits; joined waves establish visibility. Parallel Jacobi is a different preconditioner from SGS. | +| Candidate/accepted ownership and cooperative cancellation | `src/simulation.rs:649`, `src/simulation.rs:663`, `src/simulation.rs:1080`; `src/session.rs:234` | Consistent. `Simulation` commits numerical state; session quotas wrap it. No hard wall-time preemption guarantee follows from callbacks. | +| Pressure provider is a synchronous host-slice candidate boundary | `src/solver_provider.rs`; `src/provider_step.rs`; `src/recording/provider/runtime.rs` | Consistent. Pressure-only transfers and selected/fallback provenance must be disclosed; this API is not already a resident whole-step backend. | +| Particle preview is a separate physical owner and display representation | `src/particle_fluid.rs:353`; `src/material_session.rs:49`; `src/bin/desktop/history.rs` | Consistent. Diagnostic particle pressure is pre-integration; material capability says `grid_pressure_provider=false`; appearance voxels are not MAC samples. | +| Accepted interval pressure versus restart and archive failure behavior | `src/recording/carrier.rs:1`, `src/recording/carrier.rs:46`; `src/simulation.rs:312`; `src/recording.rs:320`; `docs/api/recording.md:129` | Consistent. Recording profiles have bounded admission, pressure stamps invalidate on later attempts, append failure poisons the writer, and inspection is not complete continuation state. | + +The limited slab/column, fraction and aligned-obstacle owners remain explicitly +separate from the filled-box oracle. The chapters do not promote collision +queries or particle contact to arbitrary embedded-grid pressure/stress +boundaries. Existing physical qualification failures remain historical evidence, +not new execution or a GPU pass. + +## Independently reopened primary sources + +The reviewer reopened the following sources through read-only web retrieval on +2026-10-21 instead of relying solely on the bibliography's assertion of access. +Each supports the limited claim listed; none establishes installed capabilities. + +| Primary source independently read | Claim checked and result | +| --- | --- | +| [NVIDIA asynchronous execution](https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/asynchronous-execution.html) | Stream sequencing, event dependencies, asynchronous host completion and hardware-dependent concurrency support the proposed schedule. Footer independently displayed 2026-09-10. | +| [NVIDIA graphs](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cuda-graphs.html) and [compute capabilities](https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html) | Definition/instantiation/execution and target-dependent capabilities support explicit graph/resource and admission identities. They supply no Rheon timing. | +| [NVIDIA CUDA platform](https://docs.nvidia.com/cuda/cuda-programming-guide/01-introduction/cuda-platform.html) | NVIDIA hardware, driver/toolkit, PTX and cubin/fatbin distinctions support Chapter 38's supported-device discussion. No specific model/OS certification is inferred. | +| [Khronos documentation index](https://docs.vulkan.org/spec/latest/index.html) | Independently matched Vulkan 1.4.365, generation time 2026-10-02 02:01:57Z and commit `8c9361ba8180c1f4164c0bf79de2f6e817770b0d`. This is a moving document identity, not the chosen minimum API. | +| [Khronos features](https://docs.vulkan.org/refpages/latest/refpages/source/VkPhysicalDeviceFeatures.html), [floating atomics](https://docs.vulkan.org/refpages/latest/refpages/source/VkPhysicalDeviceShaderAtomicFloatFeaturesEXT.html) and [shaders](https://docs.vulkan.org/spec/latest/chapters/shaders.html) | `shaderFloat64`, distinct buffer/shared f32/f64 atomic-add features, workgroup scope and implementation-dependent subgroup arithmetic order support capability-conditioned arithmetic and hierarchical reductions. | +| [Khronos synchronization](https://docs.vulkan.org/spec/latest/chapters/synchronization.html) | Execution dependencies alone do not supply visibility; explicit memory dependencies, host completion, semaphore and ownership rules support the resource ledger. | +| [Harris GPU Gems chapter](https://developer.nvidia.com/gpugems/gpugems/part-vi-beyond-triangles/chapter-38-fast-fluid-dynamics-simulation-gpu) and [Crane et al. GPU Gems 3 chapter](https://developer.nvidia.com/gpugems/gpugems3/part-v-physics-simulation/chapter-30-real-time-simulation-and-rendering-3d-fluids) | Confirmed historical 2D rectangular/no-free-surface scope and the 3D chapter's deliberately collocated exposition/MAC caution. Their examples are not Rheon's discretization. | +| [Bolz primary PDF](https://www.cs.columbia.edu/cg/pdfs/28_GPUSim.pdf) and [McAdams primary PDF](https://www.math.ucdavis.edu/~jteran/papers/MST10.pdf) | Confirmed GPU sparse CG/regular multigrid versus shared-memory CPU multigrid-PCG. The drafts correctly avoid transferring published timings to Rheon. | +| [Harris reduction PDF](https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf) and [Bridson/Müller-Fischer notes](https://www.cs.ubc.ca/~rbridson/fluidsimulation/fluids_notes.pdf) | Confirmed hierarchical reduction motivation, historical hardware limits and notes' title/date/preconditioner scope. Modern synchronization is separately sourced. | +| [Weber author manuscript](https://www.academia.edu/14777152/A_Cut_Cell_Geometric_Multigrid_Poisson_Solver_for_Fluid_Simulation) and [Aanjaneya author abstract](https://orionquest.github.io/papers/EGMVL/paper.html) | Confirmed geometry-specific staggered cut-cell hierarchy motivation and coupled-viscosity box smoothing. Their general suitability and detailed implementation are not inferred. | +| [Pinned AmgX PCG source](https://raw.githubusercontent.com/NVIDIA/AMGX/91a8413ef267b1c32aff4014c02820e1c5897ac2/src/solvers/pcg_solver.cu), [version file](https://raw.githubusercontent.com/NVIDIA/AMGX/91a8413ef267b1c32aff4014c02820e1c5897ac2/ReleaseVersion.txt) and [license](https://raw.githubusercontent.com/NVIDIA/AMGX/91a8413ef267b1c32aff4014c02820e1c5897ac2/LICENSES/BSD-3-Clause.txt) | Confirmed apply/dot/update/precondition structure, version 2.5.0 and BSD-3-Clause scope at the cited hash. No claim about bundled dependencies' licenses or Vulkan support follows. | +| [Pinned Flow pressure orchestration](https://raw.githubusercontent.com/NVIDIA-Omniverse/PhysX/b97365c6ddb32fda8469bc788ec18024862b0368/flow/source/nvflow/Pressure.cpp), [divergence shader](https://raw.githubusercontent.com/NVIDIA-Omniverse/PhysX/b97365c6ddb32fda8469bc788ec18024862b0368/flow/source/nvflow/shaders/PressureDivergenceCS.hlsl) and [license](https://raw.githubusercontent.com/NVIDIA-Omniverse/PhysX/b97365c6ddb32fda8469bc788ec18024862b0368/LICENSE.md) | Confirmed separate pressure passes/resources, centered collocated float4 velocity differences and BSD-3-Clause at the cited hash. No integrated MAC parity claim or code incorporation is made. | + +Independent retrieval limits are explicit: direct ACM opening returned 403; +the Bolz institutional PDF supplied an alternative primary copy. Initial +Wisconsin PDF guesses were unavailable; the cited UCDavis author PDF succeeded. +The Wiley full-text fetch timed out; the author manuscript supported only the +limited cut-cell claim. GitHub API commit metadata was inaccessible in this +reviewer session, and the PhysX commit page exceeded the retrieval size limit; +pinned raw source/license content succeeded. Exact upstream commit timestamps +remain the bibliography author's metadata check rather than an independently +reproduced reviewer check. No denied operation was retried through an access +workaround, and no source-fetch failure was treated as numerical evidence. + +## Document checks and remaining qualification + +After all repairs and renumbering, the local relative-link existence scan +covered the seven new GPU book documents and **159 local relative links, with +zero missing destinations**. No old GPU chapter paths remained in those files. +It does not validate fragment anchors, rendered URLs, equations, tables, public +downloads or PDF/Word/browser delivery. Parent-owned formatting checks and +standards review are separate evidence, not execution by this reviewer. + +No solver, numerical fixture, device enumeration, scientific workload, GPU +kernel, test suite, full book build, installation, security change or publication +was run by this reviewer. No production/shared navigation/TOC source was edited. +Passing this source review can support a documentation handoff only. Numerical +parity, device/error/cancellation qualification, end-to-end speed, memory peaks, +physical convergence and restart equivalence remain future gates in Chapter 39. diff --git a/docs/plans/gpu-grid-research/reports/standards-and-scope.md b/docs/plans/gpu-grid-research/reports/standards-and-scope.md new file mode 100644 index 00000000..e647cb0a --- /dev/null +++ b/docs/plans/gpu-grid-research/reports/standards-and-scope.md @@ -0,0 +1,34 @@ +# Standards and source-scope review + +Inspected the fetched current [Coding-Standards authority](https://github.com/MrScripty/Coding-Standards/tree/0fc92a24279f3a4574c80af0a2f31ea347a21597) +on 2026-10-21. No ancestor/workspace/Rheon `AGENTS.md`, repository `.agents/skills` +or application standards-engine tools were found. The external standards +authoring skill explicitly serves standards authors, not ordinary application +agents; it was inspected for applicability and not invoked. This is source-guided +review, not a standards-engine conformance certificate. + +Core/router select Planning/Discovery, Documentation, Commit/isolation, +Verification and Development Proportionality with its required Implementation +closure. Architecture/Code Design, Contracts, Concurrency/Resilience, Dependencies, +Licensing and Performance inform the proposed backend/lifecycle/benchmark design. +No Rust/Cargo, native FFI, GUI, persistence implementation, deployment or release +mechanism changes. Those routes must be selected against actual future code. + +| Concern | Concrete research consequence | +| --- | --- | +| Source authority | Fetch main, freeze commit and inventory actual owners; distinguish MAC pressure from particle preview voxels and historical/failed evidence | +| Planning and scope | New GPU write set, exactly one next review slice, proposed stages not production admission; disjoint authors and parent integration | +| Architecture | Independent CUDA/Vulkan allocations/build/execution; one physical/accepted-state owner; composed-design probe in plan and Chapter 39 | +| Contracts | Explicit field location/units/time/gauge/arithmetic policy, supported domains, typed failure and no silent backend fallback | +| Concurrency and resilience | Candidate transaction, bounded batches, generation-stamped copies, resource completion before reuse and durable-prefix reconciliation | +| Verification | Independent source/technical review plus document checks; no numerical, hardware or rendering evidence inferred from prose | +| Performance | Equal-tolerance identified CPU oracle; timing scopes include upload, acceptance, recording and presentation separately; no speedup promise | +| Dependencies and licensing | Verified primary sources and upstream source/license/revision; no third-party code/figures/datasets imported; future exact dependency audit | +| Documentation and credit | Existing foundations linked; only gaps expanded; Puma credit preserved; shared TOC/build/reference edits left to parent | +| Commit/lifecycle | Conventional scoped local commit after review; original checkout untouched; isolated review resource has explicit parent-owned retirement condition | + +The [source/output policy](../../../REPOSITORY_OUTPUTS.md) excludes numerical +outputs, renders, logs and archives from commits. New tracked files are authored +Markdown. The inspected educational build invokes/depends on scientific +qualification. This task executes only nonnumerical source/link/format checks; +it does not bypass build or numerical gates or relabel their old evidence fresh. diff --git a/docs/research-book/chapters/37-gpu-grid-algorithms.md b/docs/research-book/chapters/37-gpu-grid-algorithms.md new file mode 100644 index 00000000..bd37f854 --- /dev/null +++ b/docs/research-book/chapters/37-gpu-grid-algorithms.md @@ -0,0 +1,356 @@ +# 37 GPU algorithms for a compatible grid fluid + +*Research chapter for the Puma-credited Rheon book. Reviewed against the CPU source at `fd7c0f0` and primary literature on 21 October 2026. The designs below specify future work; no GPU solver, hardware result or general three-dimensional coupled liquid implementation is claimed.* + +A GPU helps a grid fluid when enough independent samples can use the same arithmetic on resident data. It cannot remove the elliptic pressure constraint, fix a mismatched boundary operator, or turn a fixed iteration count into a physical accuracy guarantee. The useful design question is therefore: which dependencies belong to the mathematics, which come from a particular numerical method, and which were introduced by the implementation? + +Chapters [3](03-staggered-grids.md)–[10](10-forces-and-viscosity.md) establish the discretization; [13](13-memory-and-parallelism.md) and [15](15-floating-point.md) establish storage and arithmetic cautions. This chapter develops their consequences for GPU algorithms without choosing CUDA or Vulkan. [Chapter 38](38-independent-cuda-and-vulkan.md) concerns backend mechanics; [Chapter 39](39-gpu-validation-and-roadmap.md) and the [implementation research plan](../implementation/gpu-research-plan.md) govern qualification. Source identifiers G1–G9 resolve to the [verified GPU bibliography](../sources/gpu-algorithm-bibliography.md). + +## 37.1 The reference problem and the transfer boundary + +Consider incompressible Newtonian fluid with positive density, a Cartesian marker-and-cell (MAC) grid, pressure at cell centers, and each velocity component on faces normal to its direction. A split step transports velocity, applies body forces and possibly viscosity, solves pressure, corrects velocity, and transports the selected scalar or interface representation. This order is a numerical choice: moving operations across one another changes splitting error. + +The bounded native reference is more specific. [`PressureWorkspace`](../../../src/pressure.rs) owns six cell-length `f64` arrays: pressure, right-hand side, residual, search direction, matrix product and preconditioned residual. It offers diagonal-Jacobi PCG and symmetric-Gauss–Seidel PCG. The ordinary [`Simulation`](../../../src/simulation.rs) stores face velocity and tracer as `f32`, prepares a candidate, checks pressure and independently reconstructed stored-velocity divergence, and commits an accepted generation. [`PressureOperator`](../../../src/operator.rs) owns the integrated SI coefficient convention. Its baseline is a connected sealed rectangular domain; additional slab, column and fixed flat-column modes have explicit restricted geometry contracts. The material particle owner and its voxelized appearance preview are separate from this MAC pressure state. + +The [CPU oracle inventory](../implementation/gpu-cpu-oracle-inventory.md) records the exact stage and ownership boundaries. A first GPU parity target should name that bounded problem, including boundary faces, zero pressure initialization, tolerance conjunction, cancellation and failed-step behavior. Merely producing similar smoke does not establish parity. A later obstacle-aware or free-surface solver is a new numerical qualification even when it reuses a pressure kernel. There is no inference from a height-column capability to arbitrary overhangs, droplets or moving cut-cell topology. + +Historical sources require the same care. Harris's 2004 chapter treats a two-dimensional homogeneous rectangular fluid, while Crane, Llamas and Tariq's 2007 chapter uses collocated textures for exposition and explicitly notes advantages of a staggered MAC grid. They demonstrate how grid passes map to GPUs, rather than supplying Rheon's pressure discretization [G1, G2]. McAdams, Sifakis and Teran develop an important parallel multigrid pressure method, but their reported implementation targets shared-memory CPU platforms [G4]. Their algorithmic ideas can motivate GPU research without becoming GPU performance evidence. + +## 37.2 Work, span, bandwidth and residence + +Let a pass have total arithmetic work W, a longest dependent chain of length S, and P effective concurrent arithmetic lanes. Even an ideal scheduler cannot complete it in less than a constant times `max(W/P, S)`. A cellwise source addition has large W and a short chain. A lexicographic triangular solve has large W but a long chain. Describing both as one loop over cells hides the difference. + +For a sequence of GPU passes, a useful accounting model is + +\[ +T_{\rm step}\approx T_{\rm upload}+T_{\rm download} ++\sum_{k=1}^{m}\left(L_k+ +\max\{B_k/\mathcal B_k,\;F_k/\mathcal P_k\}\right) ++T_{\rm dependencies}+T_{\rm host}. +\] + +Here `L` is submission/dispatch cost, B is transferred device-memory bytes, F is floating arithmetic, and the calligraphic quantities are achieved bandwidth and arithmetic rates for that workload. The maximum assumes ideal overlap within a pass and is an optimistic model, not a measured identity. Cache reuse, occupancy, addressing, branch behavior and contention affect the achieved quantities. Keeping the fluid resident removes repeated full-field transfer terms; it does not remove reduction dependencies or the host application's own memory needs. + +A seven-point constant-coefficient stencil performs only a modest amount of arithmetic per output. Its obvious uncached implementation reads seven scalars and writes one. At scalar width s that is `8sN` requested bytes for N cells. Neighboring rows reuse values, so actual external-memory traffic can be lower; coefficient loads, masks and misses can make it higher than a simple idealized model. A GPU with a large nominal arithmetic rate may spend this pass waiting on data. Conversely, adding expensive geometry or constitutive evaluation can move the bottleneck to arithmetic or control flow. + +The parallel break-even point depends on this whole expression. Small grids, small liquid islands, many tiny coarse levels, frequent host waits, short runs and resource contention can erase the benefit of high fine-grid throughput. Pressure changes that picture further: total cost is setup plus iteration count times iteration cost, measured at an accepted physical residual. A cheaper iteration that needs many more iterations may lose. Repeatedly reducing a grid diagnostic on the GPU and downloading one scalar is different from downloading all fields; the remaining scalar wait may still serialize the next decision. + +An Amdahl-style check prevents overclaiming. If fraction f of a measured end-to-end workload is accelerated by factor a while other work is unchanged, the ideal overall factor is `1 / ((1-f)+f/a)`, before transfer and integration overhead. The formula supplies a limit under stated assumptions, never a forecast for Rheon. A busy GPU, a fast pressure kernel and a faster accepted frame are three different observations. + +## 37.3 Preserve the physical operator before optimizing it + +Use the incidence convention from Chapter 3: an oriented internal face column of B has −1 at its tail and +1 at its head. Let S contain open areas, L contain normal pressure-sample distances, R contain face density, V contain liquid cell volumes, and Q contain known outward prescribed-boundary volume flux. With normal face speed u, + +\[ +D(u)=V^{-1}(Q-BSu),\qquad +G=L^{-1}B^T,\qquad +u^+=u^*-\Delta t R^{-1}Gp. +\] + +The symmetric integrated pressure system is + +\[ +Kp=b,\qquad K=BSR^{-1}L^{-1}B^T, +\qquad b=(BSu^*-Q)/\Delta t. +\] + +This assumes a two-point normal gradient on the stated orthogonal pressure-sample geometry. It does not justify arbitrary skew cut-cell centroid connections. A face transmissibility `t_e=S_e/(rho_e ell_e)` is nonnegative, shared by its two neighboring rows, and has units m⁴/kg. Cellwise gathering evaluates + +\[ +(Kp)_i=\sum_{j\leftrightarrow i}t_{ij}(p_i-p_j) ++\sum_{e\in\Gamma_D(i)}t_e p_i. +\] + +Known Dirichlet pressure contributes `t_e p_D` to b; prescribed solid speed contributes known flux. Classifying a solid face as zero-pressure air creates a different, leaking problem. + +One GPU invocation per pressure cell can gather neighboring values and write exactly one output. One invocation per correctable face can gather its two pressures and write exactly one velocity. These ownership rules avoid atomic accumulation. Face scattering can instead preserve paired contributions explicitly, but then two faces may write one cell and require a declared accumulation method. A matrix-free gather still needs one authoritative face coefficient; deriving separate copies with different clipping or geometry versions can break symmetry. + +For a closed component, constants lie in the kernel. Gauge elimination must remove a row and column consistently, or a specified mean-zero method must maintain the subspace through every operation. A pressure opening anchors its component. Closed components require compatible integrated flux before solving. A parallel connected-component procedure may converge over several passes; an unfinished label propagation is not a valid component map. Deterministic representative selection helps replay, but the correctness requirement is complete membership and correct anchoring. + +Pressure stopping must retain the physical relation + +\[ +r=b-Kp,\qquad D(u^+)=-\Delta t V^{-1}r. +\] + +If the solve uses eliminated gauges, reconstruct r on all active cells, including the removed ones. The prediction is exact for the specified algebra before rounding. It does not include the final `f32` velocity storage error. Recompute divergence from the actual stored corrected faces and the same boundary snapshot. The tolerances + +\[ +\|r\|_2\le\max(\tau_{\rm abs},\tau_{\rm rel}\|b\|_2), +\quad \|\Delta t V^{-1}r\|_\infty\le\tau_{\rm div}, +\quad \|D(u_{\rm stored}^+)\|_\infty\le\tau_{\rm stored} +\] + +are separate requirements with different units and purposes. Tiny cells amplify integrated residual; a small global norm can hide one bad local cell. + +## 37.4 A dependency graph for one candidate step + +The following graph names dependencies, not commands in a particular API. An arrow means producer completion and visibility are required before the consumer reads that version. + +~~~mermaid +flowchart TD + A[Accepted fields and geometry snapshot] --> B[Time-step diagnostics] + B --> C[Velocity backtrace and gather] + C --> D[Forces and optional viscous solve] + D --> E[Boundary flux and pressure RHS] + E --> F[Pressure iterations and reductions] + F --> G[Face pressure correction] + G --> H[Stored-velocity divergence and Courant checks] + H --> I[Tracer or interface transport] + I --> J[Finite, bound, resource and cancellation gates] + J --> K[Commit accepted generation] +~~~ + +Branches can overlap only if they read immutable versions and write disjoint resources. For example, a scalar source independent of pressure may be prepared separately, but transport driven by the projected velocity waits for that velocity. An interface update that changes pressure coefficients cannot be moved before the current pressure solve without defining which interface geometry belongs to the interval. + +Within a transport pass, output samples are independent if all sample the same immutable old velocity and source field. Updating x velocity in place and then tracing y velocity through that changed x component violates this condition. Within pressure, `Kd` is independent across rows, but the scalar `d^T Kd` must complete before alpha, and alpha must be visible before updating p and r. No API can eliminate that mathematical dependency while retaining the same recurrence. + +A workgroup-local barrier only orders participating local invocations. It does not make arbitrary other groups complete. A multi-pass phase boundary, with the required backend memory visibility, is the conservative algorithmic primitive. Spinning until another unscheduled group produces a value can deadlock. More advanced cooperative or persistent scheduling requires its own supported execution assumptions; it is not a portable baseline. Harris's reduction tutorial provides the historical tree-reduction construction and kernel decomposition motivation [G5]. Its old hardware-specific synchronization statements are not a modern API specification. + +The accepted state should remain immutable until all gates succeed. Failed pressure, invalid geometry, cancellation or insufficient sparse capacity discards the candidate. Device completion and numerical acceptance are distinct: completion makes a result available; the acceptance logic decides whether it becomes simulation state. Chapter 39 defines the evidence needed to test that lifecycle. + +## 37.5 Stencils, tiling and reductions + +### Matrix-free pressure application + +A minimal stencil contract is: + +~~~text +apply_K(snapshot, x, y): + require x and y do not alias + parallel for each active pressure cell i: + total = 0 + for each face in fixed axis/side order: + read its one authoritative coefficient and boundary class + fluid neighbor: total += t * (x[i] - x[neighbor]) + pressure boundary: total += t * x[i] + prescribed-flux wall: add no pressure coupling + y[i] = total + publish all y writes before a consumer reads them +~~~ + +The real implementation must additionally apply its reduced-system map and gauge convention. Inactive cells have declared outputs, rather than uninitialized values that enter a later reduction. A separate regular-interior kernel can avoid repeated boundary branches; a boundary kernel owns the exceptional rows. Their output regions must be disjoint and together cover the intended unknowns. + +Use x-contiguous invocations with an x-fastest layout as a candidate for neighboring memory accesses. A tile can load a shared pressure patch plus a one-cell halo and reuse it for several outputs. For a `b_x × b_y × b_z` tile, a full bounding-box halo contains `(b_x+2)(b_y+2)(b_z+2)` samples, though a seven-point stencil does not need every corner. Reuse competes with shared storage, registers, occupancy, redundant halos and synchronization. Direct cached global loads may win for some sizes. A tile shape is a tunable parameter backed by measurements, not a theorem about optimal layout. + +Temporal blocking of multiple Jacobi sweeps has a larger dependency cone. Performing k exact local sweeps needs a halo reaching k cells in the appropriate graph distance, or shrinking valid output regions. Reusing a one-cell halo for many sweeps freezes neighboring blocks and changes the global iteration to a block method. That can be useful if specified, but its residual history cannot be reported as ordinary global Jacobi. + +### Dot products, maxima and invalid flags + +Compute partial sums per workgroup, then reduce those partials through fixed additional stages. For dot products, the product precision and accumulation precision are separate choices. For a maximum diagnostic, reduce the maximum absolute finite value and separately OR an invalid-value flag; NaN behavior of a plain maximum must not decide validity accidentally. Compatibility sums also need the magnitude sum used to interpret cancellation. + +A fixed reduction tree can provide repeatability for pinned executable, precision, group size and device assumptions. It does not reproduce the CPU's ordered compensated sum. Atomic floating additions typically admit different orders; using them for a pressure norm changes numerical replay behavior even when every update is race-free. Stochastic ordering error is especially troublesome when a stopping decision changes iteration count. + +Fuse compatible vector updates and local reduction contributions to avoid rereading a vector. For example, updating r and accumulating local `r*r` contributions uses values already in registers. The global result still needs a reduction phase before a dependent decision. Fusing `Kd` with contributions to `d^T Kd` is possible when every row has the necessary d values. Fusing the scalar-dependent update into that same arbitrary multi-group pass is not justified without an additional global completion mechanism. + +## 37.6 Pressure methods on a GPU + +### Weighted Jacobi: a baseline with a spectral limitation + +For diagonal D of K, + +\[ +p^{k+1}=p^k+\omega D^{-1}(b-Kp^k). +\] + +Use distinct input and output pressure arrays per sweep. On the gauge-fixed subspace, convergence requires the eigenvalues of the error propagator `I-omega D^-1 K` to lie inside the unit disk. With K SPD and D positive, a sufficient and necessary interval for positive omega is `01` | Approximate partial counts; include every simultaneous diagnostic and scalar status | +| Corrected-advection intermediate | At least one extra transported field | Forward field survives reverse gather; reverse/correction output may share only with justified lifetimes | +| Multigrid levels | Sum of actual per-level resources | `8N/7` is only a regular cell-count estimate | +| Sparse table, tile storage and halos | Actual capacities, not active payload only | Include reserved growth and candidate topology | +| Staging/readback/history/rendering | Application-specific | Coexist with simulation and other host GPU workloads | + +For one cubic field layout matching the first three rows, with `s_u=s_q=4` and `s_p=8`, the payload is `8F+56N` bytes. Calculated examples are: + +| Cubic cells | N | F | Base payload | With one stored `f64` coefficient per face and `f64` volume per cell | +|---|---:|---:|---:|---:| +| 64³ | 262,144 | 798,720 | 20.094 MiB | 28.188 MiB | +| 128³ | 2,097,152 | 6,340,608 | 160.375 MiB | 224.750 MiB | +| 256³ | 16,777,216 | 50,528,256 | 1,281.500 MiB | 1,795.000 MiB | + +These totals deliberately differ from Chapter 13's eight-single-precision-cell-array illustration: this ledger retains the native six-double-precision pressure arrays in addition to two single-precision scalar versions. They are alternative declared layouts, not competing process-memory measurements. Reducing all pressure-work arrays to `f32` saves `24N` bytes but changes the arithmetic experiment. It cannot be called equivalent until pressure and stored-divergence gates pass on the qualified workload range. + +Resource admission must use checked products, capacity and peak overlap. A buffer fitting its individual limit does not show the full candidate step fits available device memory. A CPU fallback also needs its memory and transfer staging reserved; a fallback that tries to allocate after device failure may fail a second time. Scientific evidence and preview history have different lifetimes and should be budgeted separately. + +## 37.11 Precision, reproducibility and diagnostics + +Mixed precision is a family of designs. Separate field storage, coefficient evaluation, stencil arithmetic, local sum, global sum and scalar recurrence precision. Promoting a reduction to `f64` does not recover digits already lost in `f32` products or `Kd`. Conversely, pressure in `f64` cannot avoid rounding when a small pressure correction is finally added to a large `f32` stored velocity. + +Nondimensionalization or equation scaling can improve numerical range, provided residual-to-physical-divergence conversion retains its exact scale. Row scaling by varying cell volume usually destroys Euclidean symmetry; ordinary CG needs the integrated symmetric operator or an explicitly justified transformed inner product. A uniform scalar rescaling changes norm thresholds even when it leaves the mathematical solution unchanged. + +A first arithmetic experiment should retain the CPU's `f64` pressure work where the device supports the required operations, then isolate changes one at a time. That is a comparison design, not a requirement that every consumer device deliver useful double-precision throughput. Other candidate policies include `f32` vectors with wider reductions and true residual evaluation, or an outer higher-accuracy correction method. Each needs its own breakdown and convergence evidence. Half-width pressure storage is especially risky for high coefficient contrast and near-rest cancellation; bandwidth savings are not accuracy evidence. + +Report the gap between recursive residual and recomputed `b-Kp`. A pressure norm plateau may come from precision, poor conditioning, invalid coefficients, wrong topology, a nonsymmetric preconditioner or a stale resource. Increasing the iteration budget is informative only after those causes are separated. Curvature `d^T Kd<=0` or preconditioned product `r^Tz<=0` in an expected SPD solve is a breakdown; taking an absolute value hides a failed assumption. + +A repeatability statement should name its scope: same pinned executable/device configuration, tolerance-based cross-backend agreement, or a stronger bitwise contract actually tested. Group size, reduction tree, fused multiply-add policy, compiler transformations, subnormal handling and coefficient generation can affect results. Race freedom and determinism are separate properties; physical correctness is separate again. Bitwise repetition of a wrong boundary operator remains wrong. + +The minimal step diagnostic record should retain geometry/operator version, chosen time interval, precision policy, solver/preconditioner identity, iteration count, true full residual norms, predicted physical divergence, actual stored-face divergence, Courant value, invalid flags, termination reason and accepted generation. For sparse or multigrid runs, add active/capacity counts, hierarchy identity and setup/rebuild costs. Timings should separate pressure setup, pressure solve, transport, reductions, transfers and full accepted step. + +## 37.12 Failure mechanisms and discriminating evidence + +| Symptom | Mechanism to distinguish | Evidence that separates it | +|---|---|---| +| Pressure norm small, one cell leaks | Tiny volume or omitted gauge-cell residual | Full volume-scaled residual and stored divergence per active cell | +| GPU pressure differs but both velocities agree | Different compatible gauge | Compare pressure differences/gradients after gauge mapping | +| Smooth pressure error remains after many cheap sweeps | Weak low-frequency damping | Manufactured smooth mode and residual history versus refinement | +| Good single-pass timing, slow frame | Reductions, setup, transfer, other application work | Complete accepted-step timeline and end-to-end distributions | +| Results change with group size | Reduction order or latent race | Fixed-tree replay plus independent ownership and bounds checks | +| Grid boundary smoke smears or leaks | Sampler extension/offset/path collision | Affine offset fixture, thin-wall trace and boundary flux ledger | +| Sparse result changes with activation radius | Missing dependency support | Dense/sparse comparison with tile-edge and departing-trace fixtures | +| Multigrid fails on separated liquid pockets | Coarse topology joins components | Per-level component/nullspace maps and sealed-barrier fixture | +| SGS becomes slow despite fewer iterations | Long wavefront span or excessive joins | Preconditioner complete cost and ready-wave widths | +| Rest tank develops motion | Unbalanced force/gradient or stored precision floor | Hydrostatic cancellation, force work, stored correction and surface drift | +| Viscous material looks wrong with low divergence | Wrong stress model or inadequate temporal accuracy | Shear/channel decay and traction tests in addition to continuity | +| Cancellation corrupts next step | Candidate wrote accepted resources | Generation checks and injected failure at each phase boundary | + +Each row distinguishes a hypothesis from a diagnosis. Chapter 39 specifies experiments and evidence retention; no result in this chapter is an executed GPU fixture. Documentation/source inspection, the resource arithmetic and the stated algebraic derivations do not establish wall-clock speed, cross-device replay, a Rust-to-shader refinement or continuum validation. + +The practical sequence is to preserve the fixed-box MAC operator and acceptance lifecycle, establish independent stencil/reduction evidence, then compare pressure preconditioners at equal accepted quality. Geometry, sparse activation, constitutive complexity and coupled motion each expand the problem and should enter with their own contracts. CUDA and Vulkan can then implement a specified numerical experiment whose successes and limits are interpretable. diff --git a/docs/research-book/chapters/38-independent-cuda-and-vulkan.md b/docs/research-book/chapters/38-independent-cuda-and-vulkan.md new file mode 100644 index 00000000..24d8c914 --- /dev/null +++ b/docs/research-book/chapters/38-independent-cuda-and-vulkan.md @@ -0,0 +1,520 @@ +# 38 Independent CUDA and Vulkan compute backends + +Research synthesis, 21 October 2026. This chapter proposes two **independent +compute methods** for Puma's Rheon solver: a CUDA backend and a Vulkan compute +backend. Neither backend wraps, loads, allocates through, compiles through, or +requires the other. Both implement a shared numerical and publication contract. +All proposed interfaces and schedules below are designs, not implemented GPU +support, hardware measurements, or a claim of scientific qualification. + +Repository observations refer to main `fd7c0f0`. The +[CPU oracle inventory](../implementation/gpu-cpu-oracle-inventory.md) owns the +complete inventory; this chapter only extracts constraints that affect backend +design. [Chapter 39](39-gpu-validation-and-roadmap.md) owns prospective validation +and rollout. Official sources are identified by the `GPU-*` entries in the +[backend bibliography](../sources/gpu-backend-bibliography.md). Sources were +consulted online on 2026-10-21; moving-document revision limits are recorded +there. The live Vulkan specification identified itself as **1.4.365**, generated +2026-10-02 from commit `8c9361ba8180c1f4164c0bf79de2f6e817770b0d` [GPU-V00]. +That reference version does not set the minimum runtime API requirement. + +## A semantic interface, with two implementations + +Separate **physics**, **algorithm**, **execution backend**, and **precision** in +method identity. For example, `sealed-mac / jacobi-pcg-v1 / cuda / f32-f64-v1` +and `sealed-mac / jacobi-pcg-v1 / vulkan / f32-f64-v1` are distinct execution +methods implementing the same declared discretization. A new multigrid +preconditioner or f32-only pressure solve changes a different axis. Sharing a +method label must never hide a changed operator or tolerance. + +~~~text + numerical contract and test inputs + / \ + CUDA implementation Vulkan implementation + CUDA capability discovery Vulkan capability discovery + CUDA allocations and modules VkBuffer + VkDeviceMemory + SPIR-V + kernels, streams, events, graphs compute pipelines, queues, barriers + \ / + qualified accepted-generation record +~~~ + +This is a dependency diagram, not an API adapter stack. An optional renderer +bridge attaches to either implementation after numerical acceptance. In +particular, a CUDA-to-Vulkan rendering bridge does not turn CUDA into a Vulkan +compute backend; the Vulkan solver must execute its own SPIR-V kernels without +CUDA installed. + +The common host interface should describe inputs, ownership and completion, +without exposing either API's resource handles: + +~~~text +probe(method_requirements) -> Supported(capability_record) | Unsupported(reason) +create_session(problem, numerical_policy, memory_budget) -> Session +submit_attempt(session, accepted_generation, controls) -> AttemptTicket +poll(ticket) -> Pending | Failed(reason) | Ready(bounded_report) +commit_ready(ticket, expected_generation) -> AcceptedGeneration +request_cancel(ticket) -> cancellation_requested +export_accepted(generation, export_budget) -> snapshot_ticket +~~~ + +These are proposed operations, not additions to the current Rust API. +`Ready` means device completion and all required acceptance checks have passed; +it does not itself publish candidate arrays. Tickets carry a run/session ID, +attempt ID, old generation, method ID and buffer epoch so a stale callback +cannot publish into a replacement session. Commit checks those identities and +the latest cancellation state under the same host-side synchronization used +for publication. + +## The existing acceptance boundary constrains a GPU port + +[`simulation.rs`](../../../src/simulation.rs) stores named-axis MAC velocity and +cell tracer arrays as f32; pressure settings, time and diagnostics use f64. +`Simulation::step` promises that error or cancellation preserves accepted +fields, time and generation. Candidate buffers and pressure scratch may change. +The public `StateView` borrows host slices, so retaining this exact interface +cannot automatically yield a GPU-resident simulation. A future GPU session +needs a distinct explicit snapshot/view contract rather than silently copying +every field whenever `state()` is called. + +[`pressure.rs`](../../../src/pressure.rs) retains Jacobi-PCG and symmetric +Gauss-Seidel-PCG method identities, true residual checks and a separate physical +divergence criterion. The corrected **stored f32 velocity** has another actual +divergence gate in `simulation.rs`. A small recursive PCG residual does not +replace either true-residual qualification or the stored-velocity check. +Changing work scheduling does not authorize deleting those gates. + +[`solver_provider.rs`](../../../src/solver_provider.rs) exposes a synchronous +host-slice `PressureProvider`: its request borrows an f64 RHS and its output +fills an f64 pressure slice. Adapting a GPU pressure candidate to this existing +boundary incurs RHS upload and pressure download, followed by host acceptance. +That can be a useful bounded experiment, but it is not a full resident solver. +It must be identified separately from a port that keeps advection, RHS assembly, +pressure iterations, correction and diagnostics on the device. + +[`recording/provider/runtime.rs`](../../../src/recording/provider/runtime.rs) +explicitly states that request stamps and hashes are inspection evidence rather +than restart checkpoints or shutdown/preemption guarantees. Its accepted +context binds generation, interval, dt and prior pressure lineage. A GPU report +must preserve that distinction: completed kernel work, accepted simulation +state, recorded evidence and a durable restart file are different objects. + +The proposed shared semantic contract therefore requires the following. + +| Contract item | Required behavior in each independent backend | +| --- | --- | +| Geometry | Same named axes, staggered face counts, physical spacing, units, masks and boundary rules as the selected CPU problem | +| Operator | Same pressure coefficients, sign, volume weights, compatibility treatment and gauge; no implicit substitution of a different Laplacian | +| Precision | Explicit storage, arithmetic, accumulation and conversion policies; baseline follows f32 stored fields and f64 pressure arithmetic | +| Solver | Declared preconditioner, stopping criteria, maximum iterations and true-residual cadence; changing SGS ordering changes the method unless equivalence is established | +| Acceptance | Finite fields and scalars, valid pressure/gauge, true residual, physical divergence, stored-velocity divergence, Courant and applicable tracer/energy checks | +| Failure | Keep the old accepted generation, time and fields; invalidate candidate evidence and prevent publication | +| Publication | Advance fields, time, generation and associated report together after completion and qualification | +| Recording | Bind report and exported samples to the accepted generation and controls; do not call a pressure receipt a restart | + +The table is a proposed cross-backend contract inferred from repository +semantics. It is not a list of implemented GPU guarantees. + +## Resource residency and a budget that includes temporary storage + +Let `C = nx*ny*nz`, and let `F = Fx+Fy+Fz` be the sum of the three MAC face +counts. Use checked host arithmetic for all products, padded extents, offsets, +dispatch sizes and byte counts before any allocation. A useful lower-bound +ledger for a sealed-box prototype is below; exact scratch depends on the +selected algorithm. The CPU pressure workspace has six f64 cell arrays, which +is a starting point for budgeting, not a requirement to copy its layout. + +| Resource | Proposed residency and lifetime | Payload / additional accounting | +| --- | --- | --- | +| Accepted velocity and tracer `A` | Device; read-only throughout an attempt | `4*(F+C)` bytes | +| Candidate velocity and tracer `B` | Device; overwritten before each candidate use | Another `4*(F+C)` bytes | +| Advection/force/transport intermediates | Device; stage lifetime | Explicit method-dependent ledger; do not alias a still-read source | +| Pressure, RHS and Krylov/preconditioner scratch | Device; attempt/session lifetime | Six f64 `C` arrays would cost `48*C`; add any extra true-residual buffer | +| Geometry, masks, coefficients, optional hierarchy | Device; immutable or separately versioned | Account for materialized coefficients, halos and every multigrid level | +| Dot/norm partials and scalar state | Device; per attempt | All reduction levels, iteration counters, convergence/breakdown flags and report | +| Upload staging | Bounded host-visible storage | Initial fields, changed controls/geometry; retain until transfer completion | +| Diagnostic readback ring | Bounded host-visible storage | Fixed-size summaries plus separately requested bounded samples | +| Checkpoint/export staging | Bounded ring or tiled transfer | Reserve capacity before launch; include durable writer backlog | +| Rendering snapshot | Device; held by renderer until released | Separate buffer or explicit accepted-buffer lease; include images/mesh outputs | +| API/runtime overhead | Backend-owned host/device objects | Modules/pipelines, descriptors, command buffers/graphs, pools, alignment and allocator fragmentation | + +The budget is more than logical array payload. For Vulkan, use memory +requirements and compatible memory types, not an assumption that a buffer's +requested byte size equals its backing allocation. Host visibility, device +locality and coherence are separate properties [GPU-V06]. For CUDA, explicit +device allocations and bounded pinned staging make ownership clear; a future +allocator pool must report reserved capacity as well as live payload. Start +with allocation outside the stepping loop rather than making a pooled allocator +a correctness prerequisite. + +The normal step should upload only changed controls, leave intermediate arrays +resident, and read back a small report. Rendering and recording cadence must +not force a full field download at every pressure iteration. Unified physical +memory on some systems does not remove synchronization or ownership rules. +Oversubscription, migration, and heap exhaustion remain admission concerns; +neither backend may promise a fitting problem from grid dimensions alone. + +## CUDA: a complete native execution path + +CUDA kernels execute a grid of thread blocks and expose global/shared storage +with different lifetimes. A block barrier is local to a block [GPU-C01]. The +proposed implementation assigns contiguous x indices to neighboring threads, +uses separate component kernels when face extents differ, and bounds every +global index. First implement a straightforward global-memory stencil; shared +tiles and halo loads are later optimizations whose extra barriers must be +auditable. + +Use an explicitly created nonblocking stream for the numerical sequence. CUDA +stream order supplies dependencies within that stream. Events and +`cudaStreamWaitEvent` express dependencies across streams; host polling must +distinguish pending completion from an error. Pinned host staging supports the +intended asynchronous copies [GPU-C02]. Avoid relying on default-stream +behavior, accidental host blocking, or a device-wide synchronize per kernel. + +### CUDA kernel and dependency schedule + +~~~text +initial upload -> event(initial_ready) +compute_stream waits(initial_ready) + advect A.velocity -> velocity_tmp + add forces/boundary treatment -> B.velocity + assemble_rhs(B.velocity, geometry) -> rhs + initialize_pressure_and_krylov(rhs) -> p, r, z, d, scalar_state + repeat bounded chunks: + apply_operator(d) -> q + reduce(d*q) -> sigma; finalize alpha/breakdown + update(p,r,alpha); precondition(r) -> z + reduce(r*z and required norms) -> scalars; update d + scheduled true_residual(p,rhs) -> qualification state + all later chunk kernels honor done/breakdown state + correct_velocity(p,B.velocity) -> stored f32 B.velocity + compute actual divergence and Courant from B.velocity + transport A.tracer using the declared carrier -> B.tracer + reduce remaining acceptance diagnostics -> device_report + record event(report_ready) +readback_stream waits(report_ready) + copy fixed-size device_report -> pinned report_slot + record event(report_copied) +host observes report_copied, checks status/identity/cancellation, commits or rejects +~~~ + +The arrows denote real buffer dependencies. A reduction is at least a partial +kernel plus a final reduction kernel when more than one block participates; +`__syncthreads()` cannot finalize a grid-wide dot product. Each producer and +consumer remains in stream order or gains an explicit event dependency. A +render or checkpoint stream must similarly wait for the accepted generation +and hold a lease until its last read completes. + +A CUDA Graph can encode a stable chunk with explicit producer/consumer edges, +then be instantiated and replayed. NVIDIA separates graph definition, +instantiation and execution; graph edges constrain scheduling [GPU-C03]. This +is a launch-overhead optimization, not a new linear solver. Keep the ordinary +stream implementation as a reference path. Graph identity includes the buffer +epoch, launch extents and numerical policy; do not replay old pointers after +resize or reuse a stale captured scalar. Device graph launch and conditional +graph nodes are optional later designs, not baseline requirements. + +Chunk `B_iter` iterations at a time to limit submitted work without downloading +all dot products. Keep PCG scalars and a monotonic `done`/`breakdown` state on +the device; kernels after convergence leave solver state unchanged. A host +status copy at a chunk boundary is permitted. Choosing the chunk length trades +launch/notification overhead against cancellation latency and unnecessary +masked dispatches; it requires future measurement, not an invented constant. +True residual and maximum-iteration semantics must remain correct even when +convergence happens in the middle of a replayed chunk. + +### CUDA discovery and build independence + +Probe through CUDA's own device-management APIs: enumeration, properties, +attributes, runtime/driver compatibility and module availability [GPU-C04]. +Record device UUID, compute capability, relevant limits, driver/runtime versions +and the selected binary/PTX target. Feature support depends on the target +architecture; the official compute-capability tables are a source for mapping +requirements, not evidence about an unqueried device [GPU-C05]. FP64 support +does not imply useful FP64 throughput, and supported atomic types do not imply +deterministic accumulation. + +Proposed CUDA builds compile `.cu` kernels with the CUDA toolchain and expose a +narrow host boundary to Rust; alternatively a separately qualified driver-API +module path could load a shipped device image. NVIDIA documents the host/device +split, PTX, Cubin and fatbinary compilation [GPU-C06]. Pin compiler version, +architecture targets and arithmetic options in artifacts. A CUDA-only build +must neither invoke a SPIR-V compiler nor link a Vulkan loader. Do not claim +every installed NVIDIA device is supported by every toolkit binary. + +## Vulkan: a complete SPIR-V compute path + +Vulkan compute uses compute pipelines and workgroup invocations [GPU-V01]. A +proposed backend owns its instance, selected physical/logical device, +compute-capable queue, memory, descriptors, shader modules, pipelines and +command pools. It may share an appropriately managed Vulkan device with a +renderer, but it must also operate headlessly. Device enumeration and queue +family selection are Vulkan operations [GPU-V02]; neither calls CUDA. + +Represent numerical fields as storage buffers with fixed, documented scalar +layouts. Create buffers with storage and required transfer usage, inspect +memory requirements, choose compatible memory types, bind memory, and create +descriptor sets for the exact ranges read/written by each pipeline. Descriptor +layouts declare binding types and shader access [GPU-V03]. Use per-attempt or +per-buffer-epoch descriptor sets; do not update a set still in use unless a +separately enabled feature and its rules explicitly permit that usage. + +SPIR-V is the device program, not CUDA PTX translated at runtime. The proposed +build compiles its own shader sources for a declared Vulkan/SPIR-V target and +validates the resulting modules for that target. Record shader and compiler +hashes, specialization constants and layouts. A Vulkan-only build must neither +invoke `nvcc` nor link/load the CUDA runtime. No build scripts or dependencies +are introduced by this research chapter. + +### Supported-device differences + +The native CUDA path targets compatible NVIDIA CUDA devices and their driver, +toolkit and executable-image requirements [GPU-C09]. Vulkan selects an +implementation through its own physical-device enumeration and advertised +properties, including vendor/device identity and device type [GPU-V02]. The +Vulkan design is therefore vendor-independent at its API boundary; actual +eligibility still depends on the exact compute, memory and arithmetic profile. +An integrated GPU, discrete GPU or software implementation must be identified +as such in evidence. Merely enumerating a Vulkan device does not establish +acceleration, usable f64 throughput or numerical eligibility. + +An NVIDIA device can expose both APIs, but each backend must query its own +requirements and run its own qualification. A system with a qualifying Vulkan +implementation and no CUDA installation remains a valid Vulkan-only target. +Conversely, a headless CUDA solver need not create a Vulkan device. No specific +GPU model or operating-system support matrix is certified by this research. + +### A conservative Vulkan capability profile + +Choose a minimum API profile explicitly before implementation. One reasonable +research baseline is Vulkan 1.1 storage-buffer compute with ordinary pipeline +barriers and fences; synchronization2 is an optional, separately checked path. +This is an engineering proposal, not a claim that all Vulkan 1.1 devices run +the numerical baseline. In particular the baseline's f64 pressure arithmetic +requires `shaderFloat64`, queried and enabled explicitly [GPU-V04]. + +| Requirement / optimization | CUDA admission or selection | Vulkan admission or selection | +| --- | --- | --- | +| Compute execution | CUDA device and compatible kernel image | Physical device with `VK_QUEUE_COMPUTE_BIT`, compatible pipeline/module | +| Launch shape | Device and kernel block/shared-memory limits | `maxComputeWorkGroupCount`, `maxComputeWorkGroupSize`, `maxComputeWorkGroupInvocations`, `maxComputeSharedMemorySize` | +| Large fields | Checked allocation sizes and device budget | Storage descriptor range, allocation limits, alignment, heap/type compatibility; partition if the design supports it | +| f64 pressure | Supported target arithmetic and explicit policy; throughput measured later | Query and enable `shaderFloat64`; reject f64 method if unavailable | +| Reduction primitive | Shared-memory tree baseline; warp operations only with valid participation | Shared-memory tree baseline; subgroup arithmetic only if compute stage and operation/type support are present | +| Subgroup/warp width | Do not copy a CUDA-specific assumption into shared semantics | Query subgroup properties; never assume width 32 or 64; size control is optional | +| Floating atomic add | Exact type, address space, scope and architecture must support operation | `VK_EXT_shader_atomic_float` and the exact `shaderBufferFloat32AtomicAdd`, `shaderBufferFloat64AtomicAdd`, or shared-memory feature as appropriate | +| Float behavior | Compiler options, operation selection and target policy | Query float-control properties for bit widths and requested rounding/denormal behavior | +| Modern submission | Graph path optional | synchronization2, timeline semaphore or subgroup-size-control paths are separately gated | +| Render sharing | Optional API/device/handle compatibility | Optional renderer device/queue compatibility; external handles are not numerical prerequisites | + +Vulkan feature support must be queried and then enabled at logical-device +creation; an API version alone is not permission to use an optional feature +[GPU-V04]. Limits and subgroup properties come from properties queries +[GPU-V05]. Query extension availability before chaining extension-specific +requirements, and record both advertised and enabled values. A device lacking +f64 can report `Unsupported` for `f32-f64-v1`; a prospective mixed-precision or +f32-only method needs a distinct identity and its own qualification. It must +not silently lower arithmetic precision. + +Neither global nor shared floating atomic addition belongs in the conservative +Vulkan baseline. Their features distinguish address space, type and operation +[GPU-V07]. Use a hierarchical reduction: one workgroup writes one partial; +another dispatch reduces partials until a scalar remains. For a tail tile, +out-of-range lanes contribute zero and still reach every required workgroup +barrier. A subgroup optimization must account for active lanes and supported +types as well as `supportedOperations`; `shaderSubgroupExtendedTypes` is an +additional requirement where extended subgroup operand types need it. Vulkan +subgroup arithmetic reduction order is implementation-dependent [GPU-V08]. + +### Explicit resource dependencies and barriers + +Vulkan execution order alone does not establish memory visibility. Its +synchronization chapter distinguishes execution dependencies from availability +and visibility, and distinguishes read-after-write/write-after-write hazards +from write-after-read hazards [GPU-V09]. The following is the proposed minimum +resource dependency ledger for a single compute queue. Rows apply to actual +buffer ranges; include read/write access when the consumer modifies the range. + +| Producer -> consumer | Affected resources | Required dependency | +| --- | --- | --- | +| Upload transfer -> first compute | Initial fields, geometry, changed controls | Transfer write -> compute shader read | +| RHS assembly -> solver initialization | RHS | Compute shader write -> compute shader read | +| Operator application -> dot partials | Matrix product `q` | Compute shader write -> compute shader read | +| Reduction partials -> next reduction level | Partial buffer | Compute shader write -> compute shader read | +| Scalar finalize -> vector update | Alpha/beta, flags | Compute shader write -> compute shader read | +| Vector update -> next stencil/reduction | Pressure, residual, direction | Compute shader write -> compute shader read/write | +| True residual -> qualification/update | Residual diagnostics, status | Compute shader write -> compute shader read/write | +| Pressure -> velocity correction | Pressure | Compute shader write -> compute shader read | +| Correction -> stored-velocity diagnostics | f32 candidate velocity | Compute shader write -> compute shader read | +| Final report -> readback copy | Device report | Compute shader write -> transfer read | +| Readback transfer -> host read | Host-visible report staging | Transfer write -> host read; completion wait, then invalidate if noncoherent | +| Renderer/export reader -> next overwrite | Leased accepted/snapshot range | Reader completion before buffer reuse; not merely numerical completion | + +For the synchronization2 path, a representative same-queue buffer dependency +is the following conceptual API sequence. It omits initialization boilerplate, +not the dependency. The backend must enable `synchronization2` and use a core +or extension entry point appropriate to the selected runtime profile. + +~~~text +dispatch apply_operator(d, q) +vkCmdPipelineBarrier2( + buffer = q, offset = q_offset, size = q_bytes, + srcStage = COMPUTE_SHADER, srcAccess = SHADER_WRITE, + dstStage = COMPUTE_SHADER, dstAccess = SHADER_READ, + srcQueueFamily = IGNORED, dstQueueFamily = IGNORED) +dispatch dot_partials(d, q, partials) +barrier(partials, COMPUTE/SHADER_WRITE -> COMPUTE/SHADER_READ) +dispatch reduce_partials(partials, scalar_state) +barrier(scalar_state, COMPUTE/SHADER_WRITE -> COMPUTE/SHADER_READ) +dispatch update_vectors(...) +~~~ + +The legacy baseline expresses equivalent stages/accesses through +`vkCmdPipelineBarrier` and `VkBufferMemoryBarrier`. This is not permission to +substitute a shader `barrier()` for the cross-dispatch dependency. Workgroup +barriers only coordinate the workgroup; a shader that spins until all +workgroups arrive can deadlock when not all workgroups are resident. Use +separate dispatches for global phases. + +For separate compute and transfer queues, add semaphore dependencies with +correct wait stages. If an exclusively shared buffer crosses queue **families**, +pair release/acquire ownership transfers with the queue synchronization; +different queues in one family need synchronization but not a family transfer. +The simple first path should use one queue and avoid this extra ownership +state. Sharing mode and queue-family identities belong in every resource's +ledger. Host access must await completion, and noncoherent mapped ranges require +the specified flush/invalidate handling and atom-size alignment [GPU-V06]. + +### Vulkan submission without per-iteration host round trips + +Record a bounded command buffer containing the same numerical phases as the +CUDA schedule, with explicit barriers at the dependencies above. Ordinary +Vulkan compute does not supply a CUDA-Graph-style arbitrary device-side loop +over dispatches. A conservative design records a fixed chunk of PCG iterations +and uses device scalar flags to mask numerical updates after convergence or +breakdown. Each recorded reduction writes valid output even when masked, and +all participating invocations still obey shader barrier rules. + +Submit that chunk, copy only bounded status at its boundary, and poll a fence +or a separately supported timeline mechanism. Reuse a command buffer only +after its previous use has completed, with stable descriptor/resource epochs. +Reserve solver buffers across chunks. Re-record or select a separately recorded +variant when extents or bindings change. Dispatch-indirect control, persistent +kernels and device-generated commands are possible separate research topics; +none is needed to make Vulkan an independent backend. + +## Precision and reproducibility are part of method identity + +The CPU's compensated sums have a particular evaluation order. A GPU reduction +tree can be accurate without matching that order or its bits. Keep f64 PCG +vectors and accumulation in the first comparison profile, recompute the true +residual using the declared operator, then separately check divergence after +the f32 velocity store. An f64 pressure report alone cannot prove that the +stored field passes the physical gate. + +CUDA's floating-point appendix discusses rounding, subnormals, fused +operations and mathematical function behavior [GPU-C07]. Vulkan exposes +bit-width-specific float-control properties [GPU-V10]. A proposed method record +must state contraction/FMA, reassociation, denormal policy, rounding, division +and square-root policy, and compiler options. A compiler setting described as +fast math needs numerical qualification, not an assumption of equivalence. + +Distinguish four future claims: agreement within declared numerical bounds; +repeatability on one device/software configuration; stability across launch +shapes; and bitwise equality across backends or CPU. Evidence for one does not +establish the others. A fixed two-level reduction tree can reduce variability, +but hardware/compiler changes and subgroup reductions still need study. +Floating atomic sums introduce scheduling-dependent addition order and are a +poor default for qualification metrics. Explicitly detect nonfinite values +rather than letting a max reduction's NaN behavior hide an invalid field. + +SGS needs particular care: parallel Jacobi is pointwise, while a forward/backward +SGS sweep consumes previously updated values. Replacing lexicographic SGS with +red-black, block-Jacobi or an unsynchronized in-place sweep changes the +preconditioner. A prospective wavefront implementation must identify diagonal +dependencies and barrier phases and preserve a symmetric positive-definite +preconditioner before claiming the existing PCG method name. This is a design +constraint, not a GPU SGS equivalence result. + +## Cancellation, checkpointing and readback + +Cancellation means **stop accepting new chunks and reject the attempt**, then +retire already-submitted work before reusing its buffers. Neither an event nor +a fence is a portable command to abort an arbitrary running kernel. Avoid long +persistent kernels in the initial path. Set chunk/dispatch work bounds and +check cancellation before submission, at chunk boundaries and immediately +before commit. A user-visible latency bound additionally needs hardware/runtime +evidence; no such bound is established here. Device loss is a separate failure +class: accepted device memory may be unavailable, requiring a previously +completed durable checkpoint for recovery. + +~~~text +attempt(A_generation): + reserve B, scratch, report_slot; mark ticket unpublished + while solver not finished: + if cancel_requested: stop submitting; retire outstanding work; reject + submit one bounded chunk with explicit dependencies + poll completion/status without exposing B + submit correction, transport and all acceptance checks + await fixed-size report; reject any API error, nonfinite value or failed gate + acquire publication lock + recheck cancellation, session ID, buffer epoch and A_generation + publish B + time + generation + report atomically as one logical transaction + release publication lock + recycle old A only after all renderer/export readers release it +~~~ + +The publication operation is a host ownership change, not a giant atomic GPU +copy. Keep accepted buffers immutable while an attempt runs. Double buffering +alone is insufficient when a renderer is still reading the old accepted slot +that the next step wants as scratch; use leases/backpressure or a separately +budgeted snapshot slot. Bounded rings need an explicit policy when consumers +fall behind: skip optional visual samples, throttle exports or pause stepping +according to the requested workflow, while never publishing a partial state. + +A restart checkpoint should export a single accepted generation with geometry, +units, boundary/source controls, method and precision IDs, time/generation, +required physical state and any necessary continuation history. Solver scratch +may be omitted only when restart semantics permit reconstructing it. Record +whether pressure history is required by the selected method, rather than +treating the current interval-pressure view as a restart payload. Copy in +bounded tiles from a leased generation, record completion for every tile, +validate the manifest/checksums, then publish a complete durable file. A hash +does not prove physical correctness or successful restart. + +Readback categories should be explicit: fixed-size step report; requested +diagnostic sample; visual snapshot; and full checkpoint. Each has a byte cap, +frequency, generation stamp, completion token and backlog cap. A CPU pressure +provider adapter is an intentionally different transfer profile and must report +its full RHS/pressure transfers when evaluating performance. + +## Optional rendering interoperation + +A Vulkan solver and Vulkan renderer on one compatible device can consume an +accepted buffer or derived image using proper compute-to-graphics stage/access +dependencies, queue ownership and image layout transitions where applicable. +Rendering cannot read a candidate that may still be rejected. Copying a scalar +field into a display texture changes storage and sampling, so that texture is +not the scientific state or an acceptance oracle. + +A CUDA solver may optionally use external memory/semaphore interoperation with +a renderer. NVIDIA documents Vulkan device matching, imported memory and +signal/wait objects in its interoperability section [GPU-C08]. Matching physical +devices, supported handle types, resource lifetime and synchronization must be +qualified separately. This optional bridge neither supplies Vulkan numerical +kernels nor becomes a dependency of headless CUDA stepping. A bounded host +snapshot remains an alternative output path when no bridge is available. + +## Evidence boundary and open implementation decisions + +This chapter establishes a proposed resource/semantic design from repository +inspection and official API documents. It establishes no installed driver, +physical device, FP64 throughput, memory capacity, shader compilation, solver +accuracy, cancellation timing or speedup. No GPU workload or build was run. + +Before implementation, decide the minimum Vulkan API/SPIR-V profile, precise +f64 eligibility policy, CUDA architecture/compiler targets, shader toolchain, +Rust ownership boundary, report schema, allocation headroom and chunking policy. +Each backend needs its own failure handling and conformance evidence. Shared +acceptance fixtures and mathematical operators are desirable; shared CUDA or +Vulkan execution objects are not required by the architecture. Preserve Puma's +authorship and the existing CPU methods while qualifying each proposed method +under its own explicit identity. diff --git a/docs/research-book/chapters/39-gpu-validation-and-roadmap.md b/docs/research-book/chapters/39-gpu-validation-and-roadmap.md new file mode 100644 index 00000000..748444dd --- /dev/null +++ b/docs/research-book/chapters/39-gpu-validation-and-roadmap.md @@ -0,0 +1,355 @@ +# 39 GPU integration, validation and an implementable roadmap + +This chapter is prospective research for Puma's Rheon book, dated 21 October +2026. No GPU solver, supported-device claim, numerical qualification or timing +result is introduced. Read the [CPU source inventory](../implementation/gpu-cpu-oracle-inventory.md), +[GPU algorithms](37-gpu-grid-algorithms.md) and [independent backend paths](38-independent-cuda-and-vulkan.md) +first. The inventory freezes the inspected CPU baseline at +`fd7c0f0aee63bc8710df99f9bb01154e7619b974`; future implementation must refresh +that baseline rather than infer compatibility from this date. + +## Start with the actual oracle + +Rheon's first GPU target should be its sealed, constant-density, filled MAC box. +The initial comparison owner is `Simulation`, `PressureOperator` and +`PressureWorkspace`, including their actual f32 velocity writes and f64 pressure +arithmetic. Jacobi-preconditioned PCG is the first reference method because its +preconditioner has no triangular dependency. Symmetric Gauss-Seidel PCG remains +a second CPU cross-check. Its ordered sweeps must not become an unordered GPU +loop that is still labeled the same algorithm. + +The CPU oracle is a tuple, not the name “CPU”: source commit, toolchain/features, +grid lengths and spacing, density, initial accepted fields, source/force sequence, +time increments, solver implementation, pressure settings, corrected-divergence +and Courant gates, boundary mode, recording profile and expected failure class. +Retain the exact input bytes and grid association of each field. A future run +receipt must identify their hashes and the backend's arithmetic policy. Existing +historical audit measurements are neither fresh CPU runs nor GPU baselines. + +There are three independent comparisons. A differential check asks whether a GPU +implements the admitted discrete CPU problem. An independent analytic, rational, +dense or manufactured reference checks whether that problem and its assembly +are correct for a stated case. A refinement study checks behavior as spatial, +temporal and iterative errors are independently reduced. Agreement between two +backends can preserve a common bug; disagreement can reflect legal floating-point +order. Neither observation alone decides physical accuracy. + +The existing [pressure provider](../../../src/solver_provider.rs) is a synchronous, +host-slice candidate API for sealed boxes. It can support a later diagnostic +pressure-only prototype if its transfer cost and provenance are explicit. It is +not already an interface for a resident whole-step GPU solver. In particular, +copying host RHS and pressure every invocation does not qualify the claim that +the fluid state remains on the GPU. Adding backend names to provider metadata +cannot admit obstacle, free-surface, variable-density or viscous equations. + +## One accepted state, several execution mechanisms + +The proposed semantic contract is shared mathematics and lifecycle. CUDA and +Vulkan independently implement it with their own allocations, execution objects, +capability checks and build dependencies. A CPU-only build, a CUDA-only build and +a Vulkan-only build must remain distinct testable configurations. Installing or +initializing CUDA must not be a prerequisite for Vulkan compute. Vulkan graphics +interop is optional to CUDA computation; it cannot become a requirement merely +because the desktop renderer uses a graphics API. + +Each numerical owner continues to select its fields, associations, geometry, +boundary conditions and acceptance gates. A backend owns the representation and +execution of one explicitly admitted problem. In the current CPU path, +`Simulation` commits accepted time and generation; the session wraps that owner +with quotas and reports. A future resident numerical owner must preserve this +single commit authority; the recorder consumes accepted snapshots. The renderer consumes a +versioned projection of those snapshots. A display texture, pressure prediction +or completed dispatch is never an independent accepted-state authority. + +The dependency view is conceptual pseudocode, not a new interface declaration: + +~~~text +immutable accepted generation n + interval inputs + admitted capabilities + -> backend-owned candidate allocation + -> advection -> forces -> optional admitted viscosity + -> RHS -> pressure solve with true-residual/predicted-divergence gates + -> velocity correction -> stored-velocity divergence + interval gates + -> tracer transport/source -> complete-step gates + -> commit candidate generation n+1 + -> bounded immutable snapshot -> existing recording/render consumers +any invalid, cancelled or unsuccessful candidate -> retain generation n +~~~ + +An implementation may evaluate some reductions earlier or fuse independent work, +but it must preserve which field version a consumer reads. For example, tracer +transport using the predicted velocity instead of the accepted corrected velocity +changes the CPU algorithm. Recomputing a force at a different stage of a moving +field changes its time discretization. A faster graph with changed dependencies +is another method and needs a separate identity and qualification. + +This separation also bounds maintenance. Changing pressure preconditioning touches +the numerical backend and its method/qualification identity, not archive time +ownership. Changing an archive codec belongs to recording, not the stencil. A +Vulkan barrier repair changes execution ordering without changing the pressure +equation. Adding physical free-surface traction changes the numerical contract +and must reach every implementing backend and its independent oracles. A field +layout or arithmetic change invalidates relevant kernel and restart evidence. + +## Acceptance in physical units + +Use the same full integrated operator as the admitted CPU problem. With cell +volumes \(V_i>0\), pressure RHS \(b\), matrix \(K\), and residual +\(r=b-Kp\), the compatible pressure/correction identity is + +\[ + D u^{n+1}=-\Delta t\,V^{-1}r. +\] + +This is an exact-arithmetic identity under the operator/boundary assumptions in +Chapters 3–4. The actual code separately checks the divergence of rounded stored +velocities. In the sealed graph, \(b\) and \(Kp\) have units m³/s² and divergence +has units s⁻¹. A small Euclidean integrated residual can hide a large cellwise +divergence when a volume is small. GPU qualification must retain the full +ungauged residual, compatibility of every sealed connected component, its gauge, +and the final stored-field gate. Checking only the reduced/gauged rows misses a +physical equation. + +Define a pressure comparison only after aligning an allowed gauge. For a sealed +component one may subtract each field's own pinned-cell value, so both use the +same zero-at-pin convention. An +atmospheric or prescribed-traction boundary can fix absolute pressure; subtracting +an arbitrary mean would erase a real error. Likewise, a pressure relative norm +at a zero-pressure solution needs a stated absolute scale. + +For any scalar field \(q\), useful proposed metrics are + +\[ + E_{2,V}(q)=\sqrt{\frac{\sum_i V_i(q_i-q_i^{ref})^2}{\sum_i V_i}}, + \qquad E_\infty(q)=\max_i|q_i-q_i^{ref}|. +\] + +For face velocity, replace cell weights by the admitted face inertia or dual +volume weights. Report components and locations separately; flattening dissimilar +units into one norm has no useful physical interpretation. Compare conservative +tracer/volume updates through their boundary/source ledger; ordinary sampled +appearance tracer is not a liquid-mass conservation claim. + +An elementary mixed comparison uses + +\[ + |q_i-q_i^{ref}|\le a_q+r_q\max(|q_i|,|q_i^{ref}|). +\] + +Here \(a_q\) has q's units and \(r_q\) is dimensionless. The constants are +future, predeclared case requirements, not universal values supplied by this +chapter. They must account for discretization, iterative residual, storage and +reduction error while preserving the physical acceptance threshold. Derive them +from independent case scales and a frozen protocol before observing candidate +results. A weak CPU iteration budget is not a license to loosen a GPU gate. +When the CPU lies near a gate, compare error bounds and higher-quality oracles; +do not declare either implementation correct because it happened to pass. + +Tolerance-based parity is the normal cross-device goal. Bitwise identity is a +separate, narrower promise requiring pinned compiler/shader versions, arithmetic +options, dispatch/reduction topology and device behavior. A fixed tree can support +repeatability on a qualified configuration; it does not prove identical results +across architectures, FMA policies, subgroups or compilers. If repeated runs +produce different acceptance outcomes, record the instability and reject that +configuration until the numerical margin and ordering are understood. + +## Differential and independent checks + +The following is a future validation matrix. No row is marked executed by this +documentation task. Tests should compare intermediate stages before long-run +differences amplify and obscure the first incorrect dependency. + +| Layer | Input/problem and independent check | Required observations and negative cases | +| --- | --- | --- | +| Layout and bounds | Non-cubic grids with distinguishable x/y/z face values; analytic constant/linear sampling | Each face extent, x-fastest mapping, halo and boundary rule; tail workgroups, overflow and insufficient resource refusal | +| Operator | Tiny explicit incidence/dense matrix independent of kernel indexing | Full `Kp`, symmetry/work, nullspace, positive quadratic form on gauged subspace; closed face and disconnected-component cases only when admitted | +| Reductions | Signed dynamic-range vectors and independent high-quality accumulation | Dot, sum, max, finite flag and compatibility budget; NaN/Inf, negative curvature and reduction scratch exhaustion | +| Pressure | Fixed RHS and initial guess, analytic smooth modes and dense small solve | True full residual, gauge, iteration/failure class, pressure error and divergence identity; incompatible RHS, zero RHS and iteration limit | +| Correction | Predicted f32 MAC fields plus known pressure | Stored f32 corrected fields, exact boundary prescription, true divergence, pressure-work/kinetic ledger under its assumptions | +| Transport | Constant/linear samples, translation, midpoint backtrace and bounded source | Footprint and stage, extrema, trace error and domain escape; tracer integral is diagnostic unless flux-conservative transport is selected | +| Forces | Uniform acceleration, force density at two densities and restricted regions | Units, force sampling locations, source saturation and no write to forbidden boundary DOFs; invalid interval parameters | +| Viscosity | Existing admitted filled-box strain operator plus analytic decay/refinement | Tensor action, mass-weighted dissipation, correct dynamic/kinematic conversion and explicit stability guard; unsupported stress boundaries refused | +| Whole step | Frozen source/time sequence and actual CPU acceptance settings | Accepted time/generation, divergence, Courant, fields and diagnostics at each accepted step; every staged cancellation and rejected candidate | +| Persistence and presentation | Accepted snapshots and controlled writer failures | Exact field stages/units/identity, immutable render view, durable prefix and unrecorded tail; no failed candidate labeled accepted | +| Future geometry | Independently assembled apertures, volumes, connectivity and swept ledgers | Shared fluxes, traction, compatibility and force/work under stated model; no arbitrary mesh admission through collision-only metadata | + +The independent references in Rheon's existing +[validation material](18-validation-program.md) and +[paper audit](../appendices/G-paper-audit-and-corrections.md) have bounded scopes. +For example, the fixed-column Poiseuille solution does not validate general +three-dimensional embedded-wall viscosity. Historical failed particle rest screens +and measured cross-wall-neighborhood limitations are retained and cannot be +“fixed” by a faster grid solver. The latter were documented model limitations, +not failed preregistered barrier screens. GPU validation adds a method comparison, not a new interpretation of +old receipts or proof hypotheses. + +Refinement needs independent axes. Tighten linear tolerances until the measured +spatial/time error is insensitive to further solver tightening; vary time at +fixed sufficiently fine grid and grid at controlled time error. For an error +\(E(h)\), report the observed rate +\(p_{obs}=\log(E(h)/E(h/2))/\log 2\), the norm, reference, domain and regime. +No rate is defined by zero or noise-dominated errors. Mesh conditioning, interface +topology changes and small cut volumes require separate analysis. A bounded +field-comparison pass cannot establish general stability or convergence. + +## GPU-resident state without losing transactional publication + +The resident design keeps accepted and candidate state distinct. At the start of +a step, accepted buffers are read-only to candidate computation. On successful +completion, the session swaps or rebinds a complete candidate only after all +device writes and gates are visible. Diagnostic flags are cleared for each +attempt; device-side “converged” from an old generation is not success. + +Device loss, allocation failure, illegal execution, nonfinite arithmetic, +iteration exhaustion and user cancellation all need closed outcomes. A cancelled +kernel is generally not a transaction rollback mechanism. Stop scheduling new +batches, await or otherwise safely retire in-flight work, discard candidate +publication, and preserve accepted state. Long kernels prevent prompt cooperative +cancellation. Use bounded dispatch batches and check the cancel/generation token +between them. A device-visible flag needs the backend's valid memory and +synchronization protocol; an ordinary host store is not a portable cancellation +channel. Benchmark the resulting latency overhead as a separate tradeoff. + +Cancellation races have one owner. An observed cancellation at the final +decision checkpoint rejects the candidate; the subsequent commit has a declared +linearization point. The current CPU polls `BeforeCommit` and then commits +without another poll. A request racing after that observation may therefore +apply to the next interval; it cannot retroactively erase an accepted generation. +A future backend must declare its observation rule as precisely. A late readback from generation n must not overwrite a +display or recording result for n+1. Resource reuse requires completion tokens, +not just matching buffer dimensions. Graph/command-buffer reuse must invalidate +on geometry, descriptor, allocation or numerical-contract changes that affect +its bindings. + +Checkpoint and inspection are different capabilities. A faithful restart needs +all state that affects the next admitted step: accepted velocity/tracer or +conservative fields, time/generation, physical/configuration identity, boundary +and geometry versions, source/random state, and any history variables. If warm +pressure is restored, identify its stage and arithmetic format; if it is omitted, +declare the resulting cold start and possible convergence/bitwise differences. +PCG scratch and in-flight commands need not be restart state if checkpoints are +allowed only between complete steps. A saved image or pressure-only profile +cannot be advertised as restart. This research does not change existing recording +schemas or upgrade their capabilities. + +Readback should be scheduled by a bounded consumer policy. Small final acceptance +reports can cross each step; complete fields cross only for selected recording, +restart or inspection requests. An asynchronous staging ring overlaps copy and +later work but must retain each source generation until its copy completes. A +full ring needs an explicit choice: slow simulation, skip optional visualization, +or refuse a recording request. Silently dropping required accepted records is +not a performance optimization. GPU-resident execution cannot promise zero +transfer while a host archive requires every complete field. + +The existing accepted-state and recording failure semantics in the +[inventory](../implementation/gpu-cpu-oracle-inventory.md) remain authoritative. +An archive failure after numerical commit creates an unrecorded accepted tail; +it does not roll physics back. A future recorder adapter must reconcile uncertain +durable commits through its existing owner. Backend failure must not fabricate a +recorded acceptance, nor should a renderer fence become a physics gate unless +the admitted product contract explicitly requires it. + +## Benchmark what the user actually waits for + +GPU parallelism is useful only if useful work amortizes dispatch, synchronization, +transfers and initialization. Measure the same admitted equations to the same +physical tolerance. A low-iteration Jacobi frame and a converged CPU PCG solve +are different workloads. Iteration counts are diagnostics; a multigrid cycle, +PCG iteration and Jacobi sweep are unequal units of work. + +Predeclare workload families: small boxes exposing launch overhead; increasing +dense boxes approaching memory capacity; anisotropic dimensions; smooth and +high-frequency pressure RHS; warm and cold starts; force-rich intervals; and +only later, qualified sparse/cut geometry with controlled occupancy and topology. +Suggested size progressions are experiment inputs, not supported capacities. +Reject a configuration cleanly when its budget or descriptor/dispatch limits +cannot represent the workload. Never make an unsupported-device run part of a +geometric correctness claim. + +| Measurement | Included boundary | Interpretation | +| --- | --- | --- | +| Construction/cold start | Device discovery, allocations, shader/kernel load or compilation, pipeline/graph construction, initial upload | User startup cost, reported separately from warmed steps | +| Device mechanism | Explicit event/timestamp interval around admitted work, with dependencies complete | Stencil, reduction, solve or full-step device time; not host latency | +| Submission and synchronization | Host preparation, submissions, waits, scalar convergence/readback | Exposes roundtrip or launch limits; includes rejected attempts | +| Accepted-step latency | Input preparation through final session commit and required acceptance data | Comparable numerical service time at equal acceptance conditions | +| Recorded-step latency | Above plus required copy, serialization and durable recorder completion | Application cost when recording is mandatory | +| Presentation latency | Accepted snapshot selection through renderer-visible completion | Includes interop/staging; separate from physical step rate | +| Memory | Accepted/candidate fields, solver/preconditioner, geometry, scratch, staging and optional render views | Peak allocated/resident budgets and host mirrors reported separately | + +Use CUDA events or Vulkan timestamp queries only in their documented valid +queues and scopes; availability, timestamp precision and wrap must be checked. +A CPU wall timer around asynchronous dispatch measures submission unless it +waits for the relevant completion. Likewise, adding a global wait after every +kernel changes the performance being measured. Report both natural application +latency and scoped device timings without conflating them. + +Record exact CPU/GPU models, driver/runtime/API versions, shader/compiler flags, +available memory, power/clock policy if controlled, numerical contract, output +cadence and other material environment facts. Disclose whether host/device copies, +pipeline creation, recorder work and rendering are included. Qualify each +independent backend on its own devices; a CUDA result is not Vulkan evidence and +an NVIDIA Vulkan result is not cross-vendor evidence. + +Freeze warmup, repeated-sample and variability policy before measurement. Report +sample distributions, median and tail, observed iteration counts, accuracy and +failure rate. For successful comparable work, speedup would be +\(S=T_{CPU}/T_{GPU}\) with both timing boundaries named. No numerical value for S +is claimed here. Preserve rejected/failed attempts and capacity limits; measuring +only selected passing frames hides robustness and resource costs. Kernel tuning +must rerun relevant correctness and lifecycle gates at its new workgroup/tile +sizes, because a timing change can also expose races. + +## Proposed stages, with explicit exit gates + +These are candidate implementation milestones, not authorization to implement. +Each later stage needs a refreshed source audit, separately approved scope and +its own numerical/device evidence. CUDA and Vulkan may progress at different +rates while preserving the same contract IDs and explicit capability results. + +| Stage | Deliverable and bounded target | Gate before claiming completion | +| --- | --- | --- | +| R0 — research | This source inventory, algorithms, independent backend mappings and validation protocol | Primary-source/technical review and nonnumerical document checks; no hardware claim | +| G0 — capability/resource skeleton | Independent optional build and device discovery for each backend; no production selection | CPU-only/CUDA-only/Vulkan-only configurations, deterministic unsupported outcomes, checked dimensions and full resource budget | +| G1 — operators | Dense sealed-box gather stencils, staggered sampling, reductions and f32 correction in isolated lab | Tiny independent matrix/index references, device memory/race diagnostics when authorized, negative arithmetic/resource cases | +| G2 — pressure | Resident Jacobi PCG, full true-residual/gauge/compatibility gates | Differential pressure and corrected-field checks, independent exact/dense cases, failure/cancel gates; no full-step claim | +| G3 — bounded full step | Existing fixed-box stage order, explicit force/source, tracer and existing admitted viscosity as separately qualified choices | One-step stage comparison, long-run/refinement and transaction tests; recording adapter/profile admission separately; equal-tolerance accepted-step benchmarks | +| G4 — scaling alternatives | Symmetric parallel preconditioner, multigrid, precision and tiling candidates | Original operator preserved, convergence/accuracy/resource comparison against G2; rejected alternatives retained | +| G5 — admitted geometry families | Existing bounded static aperture/volume owner, then separately approved liquid/interface/stress families | Geometry, per-component compatibility, work/flux, traction and continuum/temporal gates for each named family | +| G6 — general solver research | Moving cut cells, coupled free-surface stress, variable density, sparse topology, capillarity/contact and later coupling | Independent mathematical/discrete validation before GPU port; swept-volume, conservation, reaction/work and interface qualifications | + +The existing canonical carrier recording profile is narrower than all of +`Simulation`'s numerical entry points: it admits its stamped `step` and +`step_with_emitter` lineage and rejects an unstamped accepted generation. A +force/viscous or GPU-resident path needs separately designed adapter/profile +qualification before claiming that recording capability. Passing existing +recording tests alone cannot admit a new numerical producer. + +G5 must not import the particle preview collision model as hydrodynamic grid +topology. G6 is not a promise that arbitrary three-dimensional liquid/solid +physics exists on CPU or GPU. Compressible conservation laws use a different +state and wave/flux contract from this incompressible projection; Chapters 26–28 +remain their research owner. Lagrangian FEM remains separately owned work. + +The first actionable future slice is G0 plus a deliberately small G1 operator +lab after implementation authorization. It must not replace runtime selection, +archive or desktop behavior. Decisions to adopt an external kernel library, +driver binding, shader compiler, sparse grid or pressure package require exact +version/license/transitive and integration review. Reading open-source examples +does not admit their code into Rheon. + +## When to stop and re-plan + +Stop the affected method when a capability is absent, an admitted operator cannot +be represented, independent tests disagree beyond their declared uncertainty, +physical acceptance changes under reduction order, or cancellation/readback +cannot preserve state ownership within its resource limit. Report exact backend, +contract, action, device and known outcome. A CPU fallback is allowed only when +the session contract explicitly authorizes it and records actual selection, +rejection and transfer cost; it cannot silently make a failed GPU benchmark pass. + +Re-plan the numerical problem when geometry changes adjointness or compatibility, +the coupled traction model needs an indefinite system, or a precision choice +cannot reach the fixed physical gate. Re-plan the execution strategy when +recording or scalar host roundtrips dominate, sparse metadata exceeds dense +storage, or the coarse multigrid solve serializes the workload. These are +observable decisions for later experiments, not reasons to promise a speedup +before one is measured. diff --git a/docs/research-book/implementation/gpu-cpu-oracle-inventory.md b/docs/research-book/implementation/gpu-cpu-oracle-inventory.md new file mode 100644 index 00000000..2103c5b5 --- /dev/null +++ b/docs/research-book/implementation/gpu-cpu-oracle-inventory.md @@ -0,0 +1,428 @@ +# CPU source inventory and oracle boundaries for GPU research + +This inventory pins the CPU reference to Rheon commit +`fd7c0f0aee63bc8710df99f9bb01154e7619b974` (`origin/main` at the start of this +research, 21 October 2026). It records source inspection, not newly executed +numerical evidence. No solver, scientific fixture, hardware benchmark, test suite, +book generator or proof checker was run for this inventory. It proposes no backend +implementation and changes no CPU algorithm, application worker or FEM code. + +The useful initial GPU oracle is the existing constant-density Cartesian MAC +carrier, with explicit per-stage and publication checks. The particle material +preview, display voxels, conservative fraction transport, restricted surface +operators and experimental Stokes solvers are different models. Their presence +does not establish one general three-dimensional liquid solver. The +[library exports](../../../src/lib.rs), [current scope audit](../../audits/general-fluid-20261019/report.md) +and [requirements roadmap](requirements-roadmap.md) identify these distinctions. +The audit's numerical observations belong to its own recorded execution and source +lineage; citing them here does not rerun or extend their qualification. + +## 1. Source map and numerical ownership + +| Owner | Actual state or operation | Reference boundary | +| --- | --- | --- | +| [geometry.rs](../../../src/geometry.rs) | Validated counts, origin, spacing, staggered indexing and managed-payload plan | Integer and physical-coordinate checks precede field allocation. | +| [simulation.rs](../../../src/simulation.rs) | Accepted/candidate velocity and tracer, physical clock, generation, step controls, final publication | Primary full sealed-box carrier oracle. Scratch and accepted state have separate ownership. | +| [sampling.rs](../../../src/sampling.rs), [advection.rs](../../../src/advection.rs) | Component-aware interpolation and midpoint backtracing of velocity/passive tracer | Bounded, clamped semi-Lagrangian transport; no conservative momentum or tracer-mass guarantee. | +| [operator.rs](../../../src/operator.rs) | Integrated pressure operator, flux RHS, correction, physical divergence | Coefficient/sign/boundary authority; changing only a pressure algorithm must preserve these equations. | +| [pressure.rs](../../../src/pressure.rs) | PCG recurrence, true residual qualification, Jacobi/SGS preconditioners | A pressure report qualifies pressure before stored f32 velocity correction. | +| [forces.rs](../../../src/forces.rs), [emitter.rs](../../../src/emitter.rs) | Explicit body force, prescribed source/nozzle controls | Acceleration and force density have distinct units; source heads do not create a solved nozzle interior. | +| [viscosity.rs](../../../src/viscosity.rs) | Explicit coupled symmetric strain on filled MAC domains | Separate opted-in workspace and stability contract; no general free-surface or imported-mesh stress boundary. | +| [parallel.rs](../../../src/parallel.rs), [scheduling.rs](../../../src/scheduling.rs), [sgs.rs](../../../src/sgs.rs) | Joined disjoint writes, adaptive CPU batching, dependent SGS sweeps | Execution decomposition; no ownership of timesteps or accepted physics. | +| [liquid_volume.rs](../../../src/liquid_volume.rs), [liquid_step.rs](../../../src/liquid_step.rs) | Cell fraction transfers and joint carrier/fraction publication | Conservative represented volume does not automatically change carrier inertia or pressure topology. | +| [free_surface.rs](../../../src/free_surface.rs), [column_surface.rs](../../../src/column_surface.rs), [column_mac.rs](../../../src/column_mac.rs) | Resolved slab/column domains and fixed flat liquid-mass bridge | Restricted surface representations, each with admission/refusal rules. | +| [static_obstacle.rs](../../../src/static_obstacle.rs), [obstacle_pressure.rs](../../../src/obstacle_pressure.rs), [obstacle_shear.rs](../../../src/obstacle_shear.rs) | Clipped single aligned box, component-aware pressure, separate reduced shear | Does not turn ordinary `Simulation` into an arbitrary triangle-mesh fluid solver. | +| [particle_fluid.rs](../../../src/particle_fluid.rs), [material.rs](../../../src/material.rs) | Particle SI state, reference-volume SPH, EOS and pair forces | Separate bounded preview with no call to the grid pressure provider. | +| [collision.rs](../../../src/collision.rs), [mesh_collision.rs](../../../src/mesh_collision.rs), [mesh_import.rs](../../../src/mesh_import.rs) | Surface queries versus particle-center static contact/import | Contact geometry and fluid pressure/shear boundary assembly are different capabilities. | +| [session.rs](../../../src/session.rs), [material_session.rs](../../../src/material_session.rs), [pressure_session.rs](../../../src/pressure_session.rs) | Run identity, immutable inspection, quotas/deadlines, synchronous step ownership | Cooperative cancellation; no hard realtime guarantee or implicit restart service. | +| [recording.rs](../../../src/recording.rs), [carrier adapter](../../../src/recording/carrier.rs), [material adapter](../../../src/recording/material.rs) | Native accepted publications in SQLite | Typed inspection/interchange with exact sample timing; incomplete integrator state. | + +## 2. Grid representation, arithmetic and memory + +For counts `(Nx, Ny, Nz)`, `GridGeometry` stores cell-centered pressure/tracer +locations and three separate face lattices: + +| Field | Shape | Physical offset from origin, in cell widths | Native scalar | +| --- | --- | --- | --- | +| x velocity | `(Nx+1, Ny, Nz)` | `(0, 1/2, 1/2)` | f32 | +| y velocity | `(Nx, Ny+1, Nz)` | `(1/2, 0, 1/2)` | f32 | +| z velocity | `(Nx, Ny, Nz+1)` | `(1/2, 1/2, 0)` | f32 | +| tracer | `(Nx, Ny, Nz)` | `(1/2, 1/2, 1/2)` | f32 | +| pressure and Krylov arrays | `(Nx, Ny, Nz)` | `(1/2, 1/2, 1/2)` | f64 | + +Flat storage is x-fast: `i + Nx*(j + Ny*k)`, with each face array's own +dimensions substituted. Singleton dimensions are supported, including a single +cell; zero counts, integer/capacity overflow, invalid positive spacing, nonfinite +origin and unrepresentable half-cell coordinates are rejected. `at_cell_density` +changes spacing, not counts. UI axis/density maxima do not guarantee allocation. +These details are explicit in [geometry.rs](../../../src/geometry.rs) and +[geometry contracts](../../../tests/geometry_contract.rs). + +`Simulation` owns two complete `Fields` sets. The candidate supplies advection +output and is later corrected in place; pressure correction needs no third +velocity set because each face reads its own prior value and immutable pressure. +`PressureWorkspace` owns six binary64 cell arrays: pressure, RHS, residual, +direction, operator product and preconditioned residual. Its SGS parallel variant +substitutes 64-bit atomic bit storage for the last array, preserving eight bytes +per value. These atomics carry stored bits, not floating-point atomic arithmetic. + +With `N=Nx*Ny*Nz` and `F` the sum of the three face counts, the baseline planned +payload is `8*F + 8*N + 48*N` bytes: paired f32 velocities, paired f32 tracer and +six f64 pressure arrays. Actual retained `Vec` capacities are checked against the +budget. This excludes allocator metadata, stack, outputs, caller objects and +process RSS. Optional boundary, viscosity, provider, fraction, recording and +rendering owners add separate allocations. A GPU estimate must additionally +account for device scratch, transfers, staging, command/runtime allocations and +simultaneously retained host copies; the CPU payload is no total-memory promise. +See [BufferPlan](../../../src/geometry.rs), [pressure allocation](../../../src/pressure.rs) +and [simulation capacity checks](../../../tests/simulation_contract.rs). + +Interpolation and physical coefficients use f64, then transport/correction +write f32 velocity/tracer. Pressure reductions use a fixed serial compensated +sum. A future change in reduction tree, fused arithmetic or storage precision can +change convergence and gate outcomes even if its mathematical expression matches. +The current CPU parallel source aims to preserve numerical order and tests exact +bits across thread counts; that is a useful CPU regression oracle, not an existing +GPU bitwise-equivalence guarantee. + +## 3. Actual accepted carrier step + +The authoritative sequence is `Simulation::prepare_impl`, followed by the +commit owner in [simulation.rs](../../../src/simulation.rs). Not every API enables +every optional stage. + +1. Invalidate the prior interval-pressure inspection stamp. Validate pause state, + positive finite requested dt, source/force/emitter inputs, workspace identity + and supported combinations. Providers are refused with a surface or prescribed + box-flux boundary on this path. Check generation/time representability. +2. Reduce dt using accepted component-speed bounds and source/external/emitter + acceleration bounds. The velocity rate is `sum_d(max|u_d|/h_d)`; active bounds + use half the configured maximum Courant allowance. Emitter motion/turbulence + can shorten dt further. The accepted report stores actual dt, which may be + shorter than the request. +3. Advect each accepted velocity component into its own staggered candidate + lattice with midpoint/RK2 backtracing and bounded trilinear sampling. Set + normal outer-box faces exactly to zero. A surface variant clears inactive + faces. The backtrace clamps both midpoint and departure to the physical box. +4. Apply legacy localized vertical smoke acceleration, then any prescribed + emitter, then explicit external body forces. Empty external force slices + omit that stage. Optional filled-domain viscosity follows force application + and precedes pressure. +5. Build/project the candidate using the ordinary pressure operator, the selected + qualified provider path, or a distinct prescribed-flux projection workspace. + Store the corrected velocity as f32. Column variants may extend one air band. +6. Recompute actual physical divergence from stored corrected velocities and + enforce `actual_divergence_limit`. Separately enforce the actual corrected + Courant bound. Neither is replaced by the pressure solver's predicted value. +7. Advect old passive tracer with corrected candidate velocity, optionally with + its explicit static tracer barrier. Apply admitted boundary tracer policy and + source/emitter tracer. Check finite values and concentration in `[0,1]`. +8. Form finite reports (including tracer integral and kinetic energy), visit the + final cancellation checkpoint, then swap candidate/accepted fields and update + clock/generation. A pressure-success stamp is attached only by the admitted + sealed-box recording entry points. + +[Sampling](../../../src/sampling.rs) clamps interpolation to donor extrema; +[advection](../../../src/advection.rs) explicitly disclaims passive-tracer mass +conservation. Midpoint trajectory evaluation does not prove second-order accuracy +of the complete split method, especially when clamping is active. Tracer's +integral is a diagnostic, not an enforced conservation residual. Kinetic energy +reported by this carrier uses its discrete face weights, not a complete +continuum energy budget. + +[Forces](../../../src/forces.rs) sample a constant vector on each component's +own face positions, skip fixed normal walls, and use half-open world-space +supports. Acceleration has units m/s²; force density N/m³ is divided by constant +rho. The work report uses the actual stored f32 velocity change and covers the +external-force stage. The legacy vertical source is prescribed acceleration, +not a coupled temperature/buoyancy equation. Source/nozzle pressure heads are +prescribed controls, not proof of a solved pipe or general traction boundary. + +## 4. Pressure equation, gauge and preconditioning + +The [operator](../../../src/operator.rs) applies the integrated symmetric matrix +`K = B diag(A/(rho*d)) B^T`. On the uniform sealed box, an interior face in +direction d has `w_d = face_area_d/(rho*h_d)`; a row sums +`w_d*(p_i-p_neighbor)` over available neighbors. Its RHS is negative outward +predicted volume flux divided by dt: + +`b_i = sum_d A_d*(u_lower_d-u_upper_d)/dt`. + +Pressure correction is +`u_new = u_star - dt*(p_positive-p_negative)/(rho*h_d)`. +Pressure is Pa; the integrated matrix product and RHS are m³/s². Coefficient +construction checks positive finite density, areas and weights. Ordinary sealed +walls require exactly zero normal velocity. Matrix application is ungauged; +PCG fixes cell zero and eliminates its search degree of freedom. + +The [pressure workspace](../../../src/pressure.rs) checks sealed RHS compatibility +using compensated sums: `abs(sum b) <= 64*epsilon_f64*sum(abs(b))`. It does not +silently subtract the mean of incompatible forcing. Every solve starts pressure +from zero; the current baseline has no pressure warm-start history. The fixed +box has one gauge, not an arbitrary obstacle-component gauge service. + +The two preserved implementations are `jacobi-pcg-v1` (default) and +`sgs-pcg-v1`. They use the same operator, pressure units, gauge, stopping criteria +and physical acceptance. Jacobi divides by the diagonal. SGS applies symmetric +forward/backward triangular solves corresponding to +`M=(D+L) D^-1 (D+L)^T` on the gauge-eliminated SPD system. It is not an +independently chosen one-way Gauss–Seidel pressure solve or multigrid method. + +For each solve, the residual norm threshold is +`max(absolute_residual, relative_residual*||b||_2)`. A recursive residual can +trigger a candidate check, but acceptance requires a freshly evaluated full +`b-Kp`, including the gauge row. Both its L2 threshold and its predicted maximum +divergence must pass. For full uniform cells, predicted divergence is +`dt*max|b-Kp|/cell_volume`; flat-column masses use their row volumes. A rejected +recursive candidate with excessive true residual replaces the residual and +restarts direction. Physical-divergence rejection alone does not gratuitously +discard conjugate progress. Nonpositive/nonfinite curvature, invalid recurrence, +iteration exhaustion and cancellation return explicit errors. + +Pressure qualification happens before f32 correction. `actual_divergence_max` +then evaluates the stored field independently. A GPU pressure comparison that +checks only pressure L2 or recursive residual misses the core acceptance contract. +Useful source oracles include the independently assembled rectangular operator, +hand chain, eliminated-gauge residual, three-dimensional manufactured pressure, +random f32 projection and retained circulation cases in +[operator contracts](../../../tests/operator_contract.rs) and +[pressure contracts](../../../tests/pressure_contract.rs). The manufactured +discrete tests are not automatically independent continuum-refinement evidence. + +The optional [provider API](../../../src/solver_provider.rs) is a candidate +generator for the sealed constant-density operator. The host reconstructs its +quality checks and can invoke reference fallback. A provider-backed accepted run +does not prove the candidate itself succeeded. Future GPU measurements must +identify attempted, rejected, selected and fallback execution and include the +appropriate transfer/qualification/correction scopes. Existing model-provider +metadata does not authorize arbitrary operator or backend substitution. + +## 5. Dependency order, cancellation and publication + +[CPU parallel helpers](../../../src/parallel.rs) split disjoint mutable scratch, +join each operation before its consumer and retain caller-thread cancellation +callbacks. Larger grid vector kernels use Rayon; small work stays serial. +Particle density/force kernels have a lower threshold because each item performs +neighbor work. [Scheduling](cpu-scheduling-contracts.md) adjusts batch sizes from +joined invocation timings. It changes decomposition, not physical timestep or +equations, and its feedback is not a repeatable speedup result. + +[SGS](../../../src/sgs.rs) has actual cross-row dependencies. Eligible large +workspaces store the preconditioned values through atomic f64 bits and traverse +tiles by `bx+by+bz`, joining every wave; the reverse triangular solve reverses +the dependency order. It calibrates 8³ and 4³ tile sweeps with scratch-only passes +and falls back for one worker or thin shapes. A GPU implementation that launches +all rows independently would change this preconditioner. Serial compensated +scalar reductions and the local per-cell arithmetic order remain separate +constraints. CPU tests compare tiled and original serial SGS bits; no GPU +execution is implied by those tests. + +`StepStage` exposes pre-advection, velocity slices, force, viscosity, pressure +iteration, projection-related, tracer and precommit stages. Availability differs +by path: ordinary in-place correction has no per-slice cancel callback, whereas +the box-flux projection has correction/acceptance callbacks. A long synchronous +kernel or one SGS preconditioner invocation cannot be asynchronously preempted +through these callbacks. Deadline checking is cooperative and has no maximum +wall latency guarantee. See [session](../../../src/session.rs), +[pressure session](../../../src/pressure_session.rs), +[simulation cancellation contracts](../../../tests/simulation_contract.rs) and +[parallel contracts](../../../tests/parallel_contract.rs). + +An error may partially overwrite scratch. It must leave accepted fields, time +and generation intact; the next attempt overwrites candidates. That guarantee +does not promise unchanged scratch, scheduler calibration, observation counters +or the previous interval-pressure stamp. The latter is invalidated on a later +attempt even if the later attempt fails. `reset` zeros fields/time and increments +generation; it is not a newly constructed run. A future GPU boundary must join +or otherwise establish completion before reads and must prevent late writes or +device errors from reaching accepted state after cancellation/publication. + +## 6. Boundaries, liquid components and collision distinctions + +| Path | Implemented contract | Limit that remains relevant to a GPU oracle | +| --- | --- | --- | +| Ordinary `Simulation` | Fixed Cartesian box, impermeable zero normal faces, constant rho | General imported-mesh fluid pressure, viscous no-slip and moving-body load coupling are absent from this step. | +| [Box flux](../../../src/box_flux.rs), [step integration](../../../src/box_flux_step.rs) | Explicit prescribed normal boundary speeds, compatibility, pressure-work and tracer-policy ledgers | A separate boundary mode, not arbitrary inlet/outlet stress selection. | +| [Slab free surface](../../../src/free_surface.rs) | Resolved wet/air interface, atmospheric zero pressure at half-cell distance | Restricted cell-aligned wet slab; not arbitrary surface topology, air inertia or capillary stress. | +| [Reconstructed columns](../../../src/column_surface.rs) | Lower-wall-connected column heights and ghost-fluid distances; unsupported scales reject | No general thin/disconnected interface, overhang or arbitrary interface momentum formulation. | +| [Liquid volume](../../../src/liquid_volume.rs) | Shared signed donor face transfers, source/inlet/outlet ledgers, outward Courant and fraction bounds | Conserves represented volume; numerical interface diffusion remains. Ordinary all-fluid carrier inertia persists even in zero-fraction cells. | +| [Filled viscosity](../../../src/viscosity.rs) | Explicit symmetric Newtonian strain, normal/shear cross-component derivatives, stationary free-slip outer box | Dynamic mu becomes nu=mu/rho; explicit bound `nu*dt*sum(1/h_d²)<=1/4`. Partial/surface/mismatched-density domains refuse on this integration path. | +| [Static obstacle](../../../src/static_obstacle.rs) | One admitted closed axis-aligned box, exact clipped cell/face measures and connected-fluid labels | Refuses unresolved cell topology. Separate pressure/reduced-shear owners do not compose a general imported mesh/liquid step. | +| [Surface query](../../../src/collision.rs), [tracer barrier](../../../src/tracer_barrier.rs) | Triangle earliest-hit/clipping and explicit tracer sampling barrier | Barrier tracer transport leaves ordinary grid velocity/pressure fixed-box physics unchanged. | +| [Particle mesh](../../../src/mesh_collision.rs) | BVH static center sweeps; two-sided open surface or connected watertight consistently wound convex solid | Contact is not an SPH pressure/shear boundary or body solver; no normal mesh motion, torque or two-way rigid-body dynamics is created. | + +Column shear/momentum and the [flat MAC bridge](../../../src/column_mac.rs), +[fitted-height geometry](../../../src/fitted_height.rs), +[translated viscous flow](../../../src/translated_viscous.rs), +[fixed-bottom ALE](../../../src/fixed_bottom_ale.rs) and +[coupled discrete flow](../../../src/coupled_discrete.rs) supply bounded research +mechanisms. Their topology, velocity families and parameter restrictions must be +carried into any comparison. Adding an extruded third velocity component does +not establish arbitrary three-dimensional geometry or flow. +The opt-in [aligned Stokes](../../../experiments/aligned_stokes.rs), +[submarine Stokes](../../../experiments/submarine_stokes.rs) and +[plate coupling](../../../experiments/plate_coupling.rs) experiments solve different +restricted problems. They are not the default carrier or a general FEM backend. + +## 7. Particle preview and display voxels + +[ParticleFluidSimulation](../../../src/particle_fluid.rs) owns accepted/candidate +`FluidParticle` arrays with f64 position, velocity, mass, temperature and material +identity, plus thermodynamic/force/diagnostic arrays and deterministic linked-cell +neighbor bins. `GridGeometry` supplies the domain and observation coordinates; +its existence does not turn the particles into a MAC pressure grid. + +The actual stages are admission and dt bounds; replayable emission into a +candidate; neighbor-bin construction; reference-volume density and EOS; +pressure/viscosity/cohesion plus gravity; stability screening; symplectic Euler +velocity then position; box-wall and optional static mesh response; diagnostic +capture and final swap. Liquid pressure uses reduced-stiffness positive +compression, gases an ideal-gas EOS; vacuum owns no particles. Density uses the +poly6 kernel and reference volumes, with an explicit 0.05 volume-density floor. +The pressure force differentiates that same poly6 compression construction; +viscosity/cohesion have their own bounded pair formulas. These details are model +choices that any particle GPU comparison must retain or explicitly requalify. + +Up to sixteen stability retries replay from accepted state with shortened dt; +failed attempts retain source credits, emission sequence, accepted particles and +clock. Density/force evaluations are disjoint per particle and read the fixed +neighbor structure, so CPU parallelism does not require unordered atomic pair +accumulation. Imported contact resolves particle-center sweeps after integration; +pair neighbor interactions are not partitioned by a thin barrier. Box-wall +clamping/normal stopping, tangential drag and heat exchange do not establish +general wall stress or resolved wetting physics. See +[material independent reference tests](../../../tests/material_reference_contract.rs) +and [mesh contracts](../../../tests/mesh_collision_contract.rs). + +The committed density, pressure and acceleration describe post-emission, +pre-integration evaluation at the old interval time. Accepted position/velocity +are at the new time, after wall response. Acceleration excludes that response. +The [material session](../../../src/material_session.rs) exposes the distinction +and explicitly reports `grid_pressure_provider=false`. Treating diagnostic +pressure as synchronized endpoint particle pressure would compare different +states. The [general-fluid audit](../../audits/general-fluid-20261019/report.md) +records failures of its particle physical-rest screens; pair/transactional +correctness does not establish continuum or experimental water accuracy. + +[PreviewFrame::capture_particles](../../../src/bin/desktop/history.rs) separately +paints emitted particles into u8 appearance voxels with a radius and max-composite +weight. Surrounding material still participates in physics even though this +preview highlights emitted material only. Carrier appearance can also be display +resampled. The [viewport](../../../src/bin/desktop/viewport.rs) extracts visible +voxel surfaces with rendering caps, and [particle view](../../../src/bin/desktop/particle_view.rs) +uses a bounded point cache. Neither creates numerical samples. Changing preview +edge, point stride, camera or surface count does not refine particles or the MAC +solver. Orange worker boxes are sampled work diagnostics, not a solver mesh or +complete utilization trace. GPU rendering speed and GPU solver speed require +separate evidence. + +## 8. Recording, inspection, restart and checkpoints + +The [accepted-publication contract](../../api/recording.md) and its source adapters +are the authority for what survives producer exit. `StepRecorder` stores one +publication and its contiguous native field BLOBs in one SQLite transaction, +with x-fast dimensions, association, scalar type, units, stage and optional +evaluation time. Canonical carrier construction has four f32 fields; an admitted +accepted sealed-box step has those plus f64 interval pressure with cell-zero +gauge. That pressure spans the actual previous time to publication time. It is +neither absolute atmospheric nor endpoint pressure. + +Canonical material recording includes accepted particle fields and their earlier +pre-integration diagnostics. Recorded source/config/control text and revisions +are provenance, not an executable stable replay codec. Mesh bytes, evaluation +positions, per-contact impulses, all source/integrator internals, pressure/Krylov +scratch and complete backend state are absent. `RecordingReader` can verify and +select a recorded publication; it cannot reconstruct complete continuation from +it. Selection must not substitute the live latest state or fabricate missing +fields. Native field bit preservation is an inspection contract, not restart +equivalence or authenticated source provenance. + +Recording follows numerical acceptance and precedes its corresponding desktop +appearance publication. Cancelled/rejected candidates add no ordinal. If physics +has accepted before cancellation, synchronous recording may finish before the +terminal cancellation report. A recording error cannot roll back accepted +physics; an unrecorded accepted tail can remain. Append failure poisons further +append admission because commit outcome may be uncertain. A fresh reader must +reconcile durable contents; acknowledged counts do not prove a failed append's +outcome. An `active` cold archive is incomplete/unknown, and destruction does +not mark a run complete. Capacity is keep-all/stop-at-limit, with no automatic +eviction, precision reduction or writer resume. See +[recording source](../../../src/recording.rs), +[carrier contracts](../../../tests/recording_carrier_contract.rs) and +[generic recording contracts](../../../tests/recording_contract.rs). + +The name `CompressibleCheckpoint` in [compressible.rs](../../../src/compressible.rs) +means a cooperative cancellation location (stage, attempt and optional cell), +not a serialized restart checkpoint. The separate compressible implementation +is a serial periodic cell-centered ideal-gas Euler/LLF solver with five conserved +averages, positivity handling and retries. Its conservative state, EOS and +boundary conditions are different from the incompressible MAC carrier and SPH +preview. GPU work must not infer a shared checkpoint/backend contract from a +type name or from all three models having three spatial coordinates. + +## 9. CPU oracle hierarchy and future validation gates + +The following are proposed gates for future authorized backend work. Existing +test source supplies reference cases; no row claims a fresh pass here. A higher +row does not make a lower row unnecessary. + +| Gate | CPU/independent oracle and required comparison | What passing would establish | +| --- | --- | --- | +| 0. Input and layout | [Geometry](../../../tests/geometry_contract.rs), [sampling](../../../tests/sampling_contract.rs): rectangular and singleton shapes, own-axis offsets, checked bounds/overflow, finite inputs | Correct indexing/representation and refusal behavior for the declared domain. | +| 1. Local operators | [Operator](../../../tests/operator_contract.rs), in-source SGS triangular hand assembly, [force](../../../tests/forces_contract.rs), [strain](../../../tests/viscosity_contract.rs) references | Signs, units, coefficients and stage equations agree with independently constructed small systems. | +| 2. Transport | [Advection](../../../tests/advection_contract.rs), clamping/donor-range and sparse-versus-dense cases | The selected transport algorithm is reproduced, including boundaries; no conservation claim follows from bounded interpolation. | +| 3. Pressure | [Pressure](../../../tests/pressure_contract.rs), both preconditioners, true full residual, gauge, compatibility, breakdown/exhaustion, actual corrected f32 divergence | Numerical pressure acceptance and its stored-field consequence match declared tolerances. Iteration count alone is insufficient. | +| 4. Stage and full step | [Simulation](../../../tests/simulation_contract.rs), [force cycle](../../../tests/forces_contract.rs), [parallel](../../../tests/parallel_contract.rs): identical accepted inputs/controls/dt with all enabled stage fields/reports | Correspondence of the implemented split carrier and acceptance order. Record CPU/GPU arithmetic differences explicitly. | +| 5. Failure lifecycle | Every reachable cancellation stage, invalid input, solver refusal and retry; [session](../../../tests/session_contract.rs) and [material session](../../../tests/material_session_contract.rs) quotas/deadlines | Accepted-state/clock/generation integrity and no delayed write after rejection. No hard realtime guarantee. | +| 6. Model-specific composition | Separate [volume](../../../tests/liquid_volume_contract.rs), [liquid step](../../../tests/liquid_step_contract.rs), [surface](../../../tests/free_surface_contract.rs), [obstacle](../../../tests/obstacle_flow_contract.rs), [material](../../../tests/material_reference_contract.rs), [mesh](../../../tests/mesh_collision_contract.rs) matrices | Only the specifically enabled physical/domain model is qualified; unsupported combinations must remain refused. | +| 7. Publication | [Carrier archive](../../../tests/recording_carrier_contract.rs), [generic archive](../../../tests/recording_contract.rs), field timing/type/association and cold selection | Inspection keeps actual accepted values with correct lineage and uncertainty semantics; no restart claim. | +| 8. Numerical/physical validation | Independent continuum manufactured solutions, spatial/temporal refinement, conservation/work ledgers and relevant experimental comparisons with uncertainty | Accuracy for the tested equations/regime, separately from CPU agreement. Existing audit findings and failures must remain visible. | +| 9. Hardware/performance | Future declared device/runtime/compiler/precision, memory peaks and complete stage/full-step timings including transfers, validation and fallback | Qualified measurements for that environment and workload, separately from correctness. No performance result currently exists in this inventory. | + +For gate 8, avoid constructing every expected value through the production +operator: CPU/GPU agreement can preserve the same assembly defect. Distinguish +an exact discrete manufactured field from a continuum source assembled +independently and refined with geometry. Exact Lean identities, regression +contracts, recorded field integrity, graphical plausibility and experimental +validation answer different questions. The existing +[audit protocol](../../audits/general-fluid-20261019/protocol.md) exemplifies +predeclared acceptance criteria and independent continuum sources; its bounded +results do not grant a general-fluid qualification. + +## 10. Book entry points and evidence limits + +The source-only review also inspected the publication entry points. The +[book README](../README.md) identifies companion experiments/depth experiments +and figure generation before the native-equation Word builder. Those scripts +are scientific execution, not passive Markdown formatting. The +[Word builder](../build_manuscript.py) explicitly selects chapter and appendix +Markdown, builds equations through Pandoc and emits DOCX/manuscript outputs; +implementation notes are not automatically in that selection. + +The [education builder](../../education/build.py) renders every chapter, appendix +and `implementation/*.md` page, but constructs its print/reading sequence using +the explicit `GUIDES` list from [native_sequence.py](../../education/native_sequence.py). +A new implementation file alone therefore supplies a web page while needing a +separate reading-order integration to enter that print sequence. The builder +validates source-bound scientific packets/reference receipts and creates portable +Markdown/source downloads. [PDF freshness](../../education/pdf_freshness.py) +binds source inputs, and [browser qualification](../../education/verify_browser.py) +renders/qualifies the PDF and interactive behavior. + +The [full research-book workflow](../../../.github/workflows/research-book-pages.yml) +executes numerical references, native contracts/examples, scientific packet +qualification, browser/Wasm work and render/publication checks; it is not a +safe substitute for a source-only documentation check under this task's scope. +The [reviewed publication workflow](../../../.github/workflows/research-book-reviewed-pages.yml) +is another publication entry point and likewise is not evidence of a local run. +No full build was attempted here. Inclusion of a note must not be described as +successful PDF/Word/browser delivery until those actual outputs and freshness +checks exist. Generated outputs remain outside Git under the +[repository output policy](../../REPOSITORY_OUTPUTS.md). + +Remaining limits are concrete: no implemented GPU oracle comparison, no device +error/cancellation qualification, no GPU memory/performance evidence, no unified +general mesh/free-surface stress integrator, no restart codec for the recorded +carrier/material path and no new physical validation from this research. These +are boundaries for subsequent authorized work, not failures concealed by a +successful source inventory. diff --git a/docs/research-book/implementation/gpu-research-plan.md b/docs/research-book/implementation/gpu-research-plan.md new file mode 100644 index 00000000..e073f191 --- /dev/null +++ b/docs/research-book/implementation/gpu-research-plan.md @@ -0,0 +1,116 @@ +# GPU grid-fluid research: reading plan and source map + +Research addition to Puma's existing book, 21 October 2026. This file is the +reader/source map; the task's [delivery plan](../../plans/gpu-grid-research/plan.md) +owns review and handoff state. The proposed implementation stages are owned by +[Chapter 39](../chapters/39-gpu-validation-and-roadmap.md), and confer no production +implementation authority. + +## Research question and source baseline + +How can Rheon's actual CPU grid-fluid algorithms execute on a GPU while +preserving their numerical and accepted-state contracts, with CUDA and Vulkan +as independently usable compute backends? The goal is a route toward a general +fluid solver whose qualified workloads can remain resident on the device. It is +not a claim that GPU execution closes currently missing fluid physics. + +The original checkout was a clean `work` branch at `6ba41c7`; it was preserved. +Remote main was fetched and inspected at +`fd7c0f0aee63bc8710df99f9bb01154e7619b974` (merged PR144). The research branch +`research/gpu-grid-algorithms-20261021` and isolated worktree +`/workspace/rheon-gpu-research` start there. Active product PR145, PR146 and the +separate Lagrangian FEM effort are outside this write set. No product source, +proof, numerical gate, shared book builder, TOC or existing bibliography is edited. + +## Exact proposed book artifacts + +| Artifact | Distinct purpose | +| --- | --- | +| [CPU oracle inventory](gpu-cpu-oracle-inventory.md) | Actual source stages, discretization, precision, boundaries, forces, validation, recording and capability limits | +| [37 GPU grid algorithms](../chapters/37-gpu-grid-algorithms.md) | Principles, mathematical dependencies and established grid-fluid GPU methods; algorithm assumptions and failures | +| [38 Independent CUDA and Vulkan](../chapters/38-independent-cuda-and-vulkan.md) | Two independent implementation paths, resource and capability contracts, synchronization and lifecycle mapping | +| [39 Validation and roadmap](../chapters/39-gpu-validation-and-roadmap.md) | CPU/independent oracles, numerical checks, benchmarks, integration and staged future implementation gates | +| [Algorithm bibliography](../sources/gpu-algorithm-bibliography.md) | Verified papers/whitepapers and pinned open-source solver source, provenance and license boundaries | +| [Backend bibliography](../sources/gpu-backend-bibliography.md) | Official NVIDIA/Khronos authority, exact consulted sections, version/access-date limits | + +Parent coordination reserves Chapters 30–36 for the separate FEM draft at +`9489aa3c1bd6652f9ab1874fce904853b9063dc3` on +`research/lagrangian-fem-book-20261021`, also based on `fd7c0f0`. +Only the new GPU chapters were moved to 37–39; existing material was not +renumbered. Parent integration owns shared navigation/coverage/build changes +and consolidation into Appendix C. The FEM proposal's book README, Chapters +11/29 and fitted-height/fitted-ALE guide edits remain untouched here. + +## Read principles before APIs + +1. Refresh the CPU inventory and identify its admitted filled-box oracle. Separate + the MAC fields from particle SPH and the appearance-density voxel renderer. +2. Read Chapter 37 after Chapters 3–5 and 13. Follow its derivation from compatible + flux operators to parallel dependency graphs, pressure methods, memory and + precision. Existing foundations are linked rather than rewritten. +3. Read Chapter 38 as two alternatives implementing those meanings. The common + contract is field/stage/acceptance semantics, not a CUDA execution object + hidden behind a Vulkan name or a Vulkan requirement on CUDA. +4. Read Chapter 39 to select future qualification evidence and implementation + slices. A paper's benchmark is not a Rheon benchmark; a solved matrix is not + a proof of PDE convergence, coupled traction or physical calibration. + +## Claim-to-source map + +| Question | Primary authority | How this research uses it | +| --- | --- | --- | +| What is already implemented? | Pinned Rheon `src/`, `tests/`, current audit and implementation notes | Source inventory, bounded domain and exact oracle; historical evidence labeled historical | +| Why do grid algorithms parallelize, and where do they synchronize? | GPU-fluid papers, GPU Gems authored technical chapters and reduction whitepaper in algorithm bibliography | Compare mechanisms and dependencies; do not transfer collocated discretization or paper timings into MAC claims | +| Which pressure strategies are candidates? | Original GPU sparse/multigrid papers, Bridson notes and pinned AMGX source | Compare Jacobi/PCG, fixed symmetric parallel preconditioners and multigrid under explicit operator assumptions | +| What does current CUDA permit? | NVIDIA programming/runtime documentation and compute-capability sections in backend bibliography | Kernel/stream/event/graph execution, synchronization, memory/resource limits and optional architecture features | +| What does current Vulkan permit? | Khronos specification, feature/property reference pages and official guide | Compute-only pipeline, SPIR-V interface, descriptor/queue/barrier/subgroup limits; no assumed FP64/atomic/subgroup size | +| How should accepted state and recordings behave? | Rheon `simulation`, `session`, `recording` and provider owners | Preserve candidate/accepted separation and durable-prefix semantics; no invented restart capability | +| Can CPU and GPU be compared meaningfully? | Actual pressure and step gates plus independently derived error metrics | Equal problem/tolerance/recording boundaries; distinguish parity, repeatability, physical convergence and performance | +| Is third-party code incorporated? | Each upstream source's pinned license, Rheon Apache-2.0 | No code, figures, binaries or datasets incorporated; future adoption needs a separate exact dependency/license audit | + +The bibliographies distinguish immutable source revisions from moving official +documentation. Access date is 2026-10-21; a moving page is evidence of consulted +content, not a reproducible SDK build or an assertion about every installed +driver. A later implementation must pin its actual compiler, API environment, +SPIR-V target, loader/binding versions and device features. Source version/date +and license are source-review facts, not hardware verification. + +## Existing material and the actual gaps + +Chapter 13 already discusses SoA layouts, reductions, CPU scheduling and memory. +Chapters 3–5 already derive staggered operators and pressure, and Chapters 15/18 +already distinguish floating arithmetic and validation. The new material adds +the missing GPU dependency/execution mapping, parallel preconditioner analysis, +capability-conditioned precision, independent CUDA/Vulkan architecture, resident +transaction/checkpoint/readback design and implementable differential benchmark +protocol. It does not duplicate a generic backend design as evidence of a solver. + +Chapters 8–10 and 19–29 own physical boundary, surface, stress and coupling +assumptions. GPU chapters explain what additional data/dependencies their admitted +variants need. They do not promote a static query mesh, reduced column, bounded +cut reference or particle collision response into general evolving 3D grid +physics. Compressible and Lagrangian FEM owners remain separate. + +## Build scope and safe review boundary + +The book is source Markdown. `docs/education/build.py` discovers chapters and +appendices, plus implementation guides; it validates reference/source bindings +and uses Pandoc and npm assets. `docs/education/generate_references.py` executes +scientific companion/expansion generators. The Pages workflow also executes +native numerical qualification, fixture generation and browser/PDF freshness +checks. The manuscript builder calls Pandoc and writes generated Word/Markdown +artifacts. These commands were inspected as source and were not executed here. + +Only nonnumerical document/link/format and source-scope checks are authorized in +this task. No GPU/hardware or scientific workload, driver/toolchain installation, +security change, full book build, render, merge, deployment or automatic +publication is part of the evidence. No generated numerical artifact or copied +research paper is committed. Existing author credit remains Puma. + +Parent integration should inspect source/PDF/browser freshness rules and the +complete workflow again before choosing any later build. The Markdown tables +and pseudocode need real output-format review: the current Word builder does +not have a dedicated Markdown-table renderer. A passing local link check alone +does not demonstrate readable Word tables, correct equations, public site links +or a new published edition. The new source bibliographies are linked by the +chapters; their appearance in standalone downloads needs parent build review. diff --git a/docs/research-book/sources/gpu-algorithm-bibliography.md b/docs/research-book/sources/gpu-algorithm-bibliography.md new file mode 100644 index 00000000..baa9a3d8 --- /dev/null +++ b/docs/research-book/sources/gpu-algorithm-bibliography.md @@ -0,0 +1,87 @@ +# GPU grid algorithm bibliography and source audit + +This is the source supplement for [Chapter 37](../chapters/37-gpu-grid-algorithms.md) of the Puma-credited Rheon research book. Access date for every successful retrieval below: **2026-10-21**. “Verified” means the identified content or source file was inspected through read-only web retrieval; it does not mean its implementation was built, executed, benchmarked, or independently numerically validated. Local design deductions in Chapter 37 are identified as such and do not acquire validation from a citation. + +The scope of each entry matters. A paper can motivate an algorithm without proving its suitability for Rheon's boundary-value problem. A published timing remains a result for the authors' implementation, machine and stopping criterion. A source repository supplies inspectable implementation evidence, with separately identified version and license; no source was incorporated into Rheon in this research change. + +## G1 — Harris: grid passes on an early GPU + +**Mark J. Harris.** “Fast Fluid Dynamics Simulation on the GPU.” *GPU Gems*, chapter 38, 2004. [Publisher-hosted chapter](https://developer.nvidia.com/gpugems/gpugems/part-vi-beyond-triangles/chapter-38-fast-fluid-dynamics-simulation-gpu). + +Inspected the full HTML chapter, especially sections 38.1.3 and 38.3. It describes a two-dimensional, homogeneous rectangular-domain stable-fluids implementation using GPU texture passes. It explicitly excludes free surfaces from that treatment. Used for the historical decomposition into independent grid operations; its old GPU timings and texture/API constraints are not current Rheon evidence. No accompanying sample software was downloaded or licensed for incorporation. + +## G2 — Crane, Llamas and Tariq: three-dimensional visual-effects pipeline + +**Keenan Crane, Ignacio Llamas and Sarah Tariq.** “Real-Time Simulation and Rendering of 3D Fluids.” *GPU Gems 3*, chapter 30, 2007. [Publisher-hosted chapter](https://developer.nvidia.com/gpugems/gpugems3/part-v-physics-simulation/chapter-30-real-time-simulation-and-rendering-3d-fluids). + +Inspected HTML sections 30.2.3 onward for grid storage, simulation passes and corrected advection. The chapter explicitly uses cell-centered quantities for exposition while identifying staggered MAC discretization as more robust. Used for GPU pipeline history and as a warning against importing a collocated sample's pressure stencil into Rheon's MAC operator. Its fluid/rendering examples do not establish conservation or boundary correctness for the proposed backend. + +## G3 — Bolz et al.: sparse CG and regular-grid multigrid + +**Jeff Bolz, Ian Farmer, Eitan Grinspun and Peter Schröder.** “Sparse Matrix Solvers on the GPU: Conjugate Gradients and Multigrid.” *ACM Transactions on Graphics* 22(3), 917–924, 2003. DOI: [10.1145/882262.882364](https://doi.org/10.1145/882262.882364). [Caltech-hosted paper PDF](https://multires.caltech.edu/pubs/GPUSim.pdf). + +Verified the eight-page primary PDF, including its solver formulation, stream/gather organization and reductions. Its GPU work includes sparse CG and regular-grid multigrid. Used to distinguish local vector/stencil work from global scalar dependencies. The text's GeForce FX execution model is historical. Direct ACM retrieval returned HTTP 403; the Caltech primary paper supplied the inspected content. The initially tried `gpusolver.pdf` address was unavailable and is not a supporting source. + +## G4 — McAdams, Sifakis and Teran: parallel multigrid preconditioning + +**Aleka McAdams, Eftychios Sifakis and Joseph Teran.** “A Parallel Multigrid Poisson Solver for Fluids Simulation on Large Grids.” *ACM SIGGRAPH/Eurographics Symposium on Computer Animation*, 2010. [Author/institution publication record](https://graphics.cs.wisc.edu/Papers/2010/MST10/). [Author-hosted full paper PDF](https://www.math.ucdavis.edu/~jteran/papers/MST10.pdf). + +Verified the ten-page PDF and publication metadata. The work combines a lightweight geometric multigrid cycle with CG for irregular voxelized pressure domains and mixed Neumann/Dirichlet conditions. Its implementation and reported measurements target **shared-memory CPU platforms**. Used for pressure algorithm design, hierarchy overhead and boundary-aware preconditioning; it is not labeled a GPU implementation or a GPU benchmark in this chapter. + +## G5 — Harris: hierarchical reductions + +**Mark Harris.** *Optimizing Parallel Reduction in CUDA*. NVIDIA Developer Technology presentation. [NVIDIA-hosted 38-page PDF](https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf). + +Verified the PDF title, author, tree reduction, partial-result decomposition and optimization discussion. A publication date is not asserted because the inspected title slide does not establish one. Used only for the historical reduction construction and bandwidth motivation. Old hardware-specific claims about global synchronization and warp execution are not used as current CUDA or Vulkan normative requirements. Backend synchronization is governed by the separately reviewed API sources in Chapter 38. + +## G6 — Bridson and Müller-Fischer: pressure, MIC and boundary derivations + +**Robert Bridson and Matthias Müller-Fischer.** *Fluid Simulation: SIGGRAPH 2007 Course Notes*. Title page dated **10 August 2007**. [University of British Columbia full PDF](https://www.cs.ubc.ca/~rbridson/fluidsimulation/fluids_notes.pdf). + +Verified the 93-page PDF, its title page and contents including sections 4.3.2–4.3.4 on CG, incomplete Cholesky and modified incomplete Cholesky, plus curved pressure boundaries. Used as the primary course-note reference for CPU-oriented preconditioning and discrete fluid semantics. A first retrieval timed out; a later full PDF retrieval succeeded. The notes are not a GPU parallel triangular-solve implementation. + +## G7 — Weber et al.: cut-cell multigrid as geometry-specific research + +**Daniel Weber, Johannes Mueller-Roemer, André Stork and Dieter Fellner.** “A Cut-Cell Geometric Multigrid Poisson Solver for Fluid Simulation.” *Computer Graphics Forum* **34(2)**, 481–491, 2015. DOI: [10.1111/cgf.12577](https://doi.org/10.1111/cgf.12577). [Publisher article](https://onlinelibrary.wiley.com/doi/full/10.1111/cgf.12577). [Author institute publication listing](https://www.igd.fraunhofer.de/en/institute/employees/daniel-weber.html). + +Verified publisher-indexed metadata and author-institute listing; also inspected the primary author-manuscript text in a [hosted manuscript copy](https://www.academia.edu/14777152/A_Cut_Cell_Geometric_Multigrid_Poisson_Solver_for_Fluid_Simulation). The limited use here is its staggered finite-volume cut-cell hierarchy motivation. No numerical rate or GPU timing is imported. Direct publisher opening timed out/returned 403; the Fraunhofer repository record was protected by an access-denial page. Those failed fetches are not represented as full-text verification. Some institutional indexing displays volume 35; the publisher citation and manuscript identify volume **34**, used here. + +## G8 — NVIDIA AmgX: inspected GPU linear-solver source + +**NVIDIA. AmgX.** Repository [NVIDIA/AMGX](https://github.com/NVIDIA/AMGX). Inspected commit **`91a8413ef267b1c32aff4014c02820e1c5897ac2`**, GitHub commit timestamp **2026-07-09T07:00:23Z**. The pinned [ReleaseVersion.txt](https://github.com/NVIDIA/AMGX/blob/91a8413ef267b1c32aff4014c02820e1c5897ac2/ReleaseVersion.txt) reads **2.5.0**. This identifies the inspected snapshot; it is not a claim that the commit is a version-release date. + +Inspected through raw source retrieval: + +- [PCG implementation](https://github.com/NVIDIA/AMGX/blob/91a8413ef267b1c32aff4014c02820e1c5897ac2/src/solvers/pcg_solver.cu): operator application, inner products, vector updates and preconditioner invocation. +- [Polynomial smoother source](https://github.com/NVIDIA/AMGX/blob/91a8413ef267b1c32aff4014c02820e1c5897ac2/src/solvers/chebyshev_poly.cu): an implementation reference, not an endorsed parameter selection. +- [Example PCG/aggregation/Jacobi configuration](https://github.com/NVIDIA/AMGX/blob/91a8413ef267b1c32aff4014c02820e1c5897ac2/src/configs/PCG_AGGREGATION_JACOBI.json): demonstrates composition; its symmetry or suitability is not inferred from the outer solver name. +- [BSD 3-Clause license text](https://github.com/NVIDIA/AMGX/blob/91a8413ef267b1c32aff4014c02820e1c5897ac2/LICENSES/BSD-3-Clause.txt), plus source SPDX header on the polynomial smoother. + +License conclusion is limited to the inspected project license/source headers; bundled dependencies have their own notices. This is a CUDA solver-library comparison, not a Vulkan solution, full fluid pipeline, or proposed mandatory Rheon dependency. No library build, installation or execution occurred. + +## G8b — NVIDIA Flow in PhysX: inspectable sparse pressure passes + +**NVIDIA. Flow SDK**, in [NVIDIA-Omniverse/PhysX](https://github.com/NVIDIA-Omniverse/PhysX). Inspected commit **`b97365c6ddb32fda8469bc788ec18024862b0368`**, GitHub commit timestamp **2026-10-15T11:12:53Z**. The commit is the reproducible version identifier; no independent Flow semantic-version number was established. + +Verified these pinned files via raw source retrieval: + +- [Pressure pass orchestration](https://github.com/NVIDIA-Omniverse/PhysX/blob/b97365c6ddb32fda8469bc788ec18024862b0368/flow/source/nvflow/Pressure.cpp): separate divergence, Jacobi, residual, restriction, prolongation and subtraction pipelines, plus pressure-level resources. +- [Jacobi shader](https://github.com/NVIDIA-Omniverse/PhysX/blob/b97365c6ddb32fda8469bc788ec18024862b0368/flow/source/nvflow/shaders/PressureJacobiCS.hlsl): six-neighbor pressure texture reads, table-based addressing and distinct output. +- [Divergence shader](https://github.com/NVIDIA-Omniverse/PhysX/blob/b97365c6ddb32fda8469bc788ec18024862b0368/flow/source/nvflow/shaders/PressureDivergenceCS.hlsl): centered differences of collocated `Texture3D` velocity components. This explicitly limits comparability with Rheon's MAC discretization. +- [Root BSD 3-Clause license](https://github.com/NVIDIA-Omniverse/PhysX/blob/b97365c6ddb32fda8469bc788ec18024862b0368/LICENSE.md), BSD-3-Clause SPDX headers on inspected pressure files, and [Flow README with third-party acknowledgments](https://github.com/NVIDIA-Omniverse/PhysX/blob/b97365c6ddb32fda8469bc788ec18024862b0368/flow/README.md). + +This modern source snapshot is distinct from the older `NVIDIAGameWorks/Flow` repository and its license. A guessed `LICENSE.md` in that older repository returned 404 and is not evidence for its license. No claim is made that the inspected shaders implement the integrated MAC pressure equation or physically general liquid coupling. No Flow binary was built or run. + +## G9 — Aanjaneya et al.: coupled viscous multigrid + +**Mridul Aanjaneya, Chengguizi Han, Ryan Goldade and Christopher Batty.** “An Efficient Geometric Multigrid Solver for Viscous Liquids.” *Proceedings of the ACM on Computer Graphics and Interactive Techniques* 2(2), article 14, 2019. DOI: [10.1145/3340255](https://doi.org/10.1145/3340255). [Primary author project page and abstract](https://orionquest.github.io/papers/EGMVL/paper.html). + +Verified the author page's authorship, venue, year and abstract describing a box smoother for coupled viscous stencils and multigrid preconditioning. The full [author-hosted PDF](https://cs.uwaterloo.ca/~c2batty/papers/Aanjaneya2019/Aanjaneya2019.pdf) exceeded the web tool's content-size limit; direct ACM opening returned 403. Therefore this entry supports the narrow abstract-level point about coupled smoothers, not a detailed reproduction of the paper or an implementation audit. No published performance figure is transferred to Rheon. + +## Retrieval and qualification limits + +All repository hashes and dates above came from read-only GitHub API metadata; source contents and licenses were fetched at those hashes. API default branches were used only to resolve a snapshot and are not the cited version identity. Project license texts are distinguished from licenses of bundled third-party material. A final implementation decision must review the exact source selected for reuse and retain its notices. + +The inspected sources span graphics-oriented stable-fluids examples, pressure linear algebra, parallel CPU multigrid and GPU solver source. Particle-only solvers, lattice-Boltzmann methods and appearance voxel grids are not treated as evidence of MAC projection compatibility. This bibliography intentionally makes no unsupported claim that a particular package offers identical independently maintained CUDA and Vulkan fluid solvers. + +The research session performed source and documentation inspection only. It did not install drivers/toolchains, use GPU hardware, execute scientific kernels, run fluid fixtures, reproduce paper benchmarks, or run the full research-book build. Remaining scientific and performance qualification belongs to [Chapter 39](../chapters/39-gpu-validation-and-roadmap.md). diff --git a/docs/research-book/sources/gpu-backend-bibliography.md b/docs/research-book/sources/gpu-backend-bibliography.md new file mode 100644 index 00000000..12f83362 --- /dev/null +++ b/docs/research-book/sources/gpu-backend-bibliography.md @@ -0,0 +1,270 @@ +# Independent GPU backends: official-source bibliography + +Accessed online **2026-10-21**. This ledger supports +[Chapter 38](../chapters/38-independent-cuda-and-vulkan.md). Primary NVIDIA and +Khronos documentation only; repository observations are separately linked at +the end. API facts support the design, not a claim that Rheon implements them +or that a particular GPU was queried or measured. Proposed schedules, +interfaces, memory ledgers and admission policies are engineering synthesis. + +## Revision policy + +NVIDIA's current **CUDA Programming Guide** pages display **Last updated on +Sep 10, 2026** where individually verified below. The guide's landing page does +not print a toolkit version in the retrieved text; this ledger does not infer +one from another document. The separately retrieved CUDA Runtime API reference +does explicitly identify **13.4** in its page title. Moving URLs are preserved +as accessed snapshots, not mislabeled immutable version URLs. + +Khronos's [documentation index](https://docs.vulkan.org/spec/latest/index.html) +explicitly identifies the consulted Vulkan Specification build as **1.4.365**, +including all registered extensions, generated **2026-10-02 02:01:57Z**, branch +`main`, commit **`8c9361ba8180c1f4164c0bf79de2f6e817770b0d`**. Specification +chapter URLs below are that moving documentation build as observed on the +access date. Standalone reference pages are also moving documents; their +retrieved pages do not separately print a version/build identifier, so no +independent revision identity is asserted for them. A future implementation +must pin and recheck its actual headers, loader, driver, compiler and target +environment rather than assuming the latest documentation is its runtime. + +## NVIDIA CUDA + +### GPU-C00. CUDA Programming Guide: identity and scope + +[Guide landing page](https://docs.nvidia.com/cuda/cuda-programming-guide/) and +[§1.1 Introduction](https://docs.nvidia.com/cuda/cuda-programming-guide/01-introduction/introduction.html). + +Verified: official guide identity; §1.1.1 describes CUDA's independence from +graphics APIs. Introduction footer explicitly dated 2026-09-10. No toolkit +version inferred. Access: public web retrieval, 2026-10-21. + +### GPU-C01. §2.3 Writing SIMT Kernels + +[Official chapter](https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html). + +Consulted: §2.3.2 Thread Hierarchy; §2.3.2.1 Thread Block Synchronization; +§2.3.3 GPU Device Memory Spaces, especially global/shared memory; §2.3.5 +Atomics. Supports block-local synchronization, memory scope distinctions and +kernel decomposition. Individually verified footer: 2026-09-10. Access: +2026-10-21. Does not establish a device's supported block size or occupancy. + +### GPU-C02. §2.5 Asynchronous Execution + +[Official chapter](https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/asynchronous-execution.html). + +Consulted: §2.5.1 asynchronous completion; §2.5.2 CUDA Streams, including +§2.5.2.3 transfers and §2.5.2.4 synchronization; events and cross-stream +dependencies; §2.5.6 default/blocking/nonblocking stream behavior. Supports the +independent CUDA scheduling design and pinned host staging. Individually +verified footer: 2026-09-10. Access: 2026-10-21. An asynchronous API is not a +guarantee of hardware concurrency or a timing bound. + +### GPU-C03. §4.2 CUDA Graphs + +[Official chapter](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cuda-graphs.html). + +Consulted: §4.2.1 Graph Structure; §4.2.2 Building and Running Graphs; +§4.2.3 Updating Instantiated Graphs; §4.2.6 Device Graph Launch; +§4.2.7 Using Graph APIs. Supports graph dependency edges, separate definition, +instantiation and execution, and why advanced graph paths need distinct +requirements. Individually verified footer: 2026-09-10. Access: 2026-10-21. +Chapter 38's bounded PCG chunk is a proposed design, not vendor sample code. + +### GPU-C04. CUDA Runtime API Reference Manual 13.4, §6.3 Device Management + +[Official device-management reference](https://docs.nvidia.com/cuda/cuda-runtime-api/cuda_runtime_api/group__CUDART__DEVICE.html). + +Verified page title: **CUDA Runtime API Reference Manual 13.4 documentation**. +Consulted function inventory and descriptions for device enumeration, +`cudaGetDeviceProperties` and `cudaDeviceGetAttribute`. Supports CUDA-owned +capability discovery. No last-updated date was displayed in retrieved text; +access date is 2026-10-21. This version applies to this reference, not an +inferred installed toolkit or Programming Guide version. + +### GPU-C05. §5.1 Compute Capabilities + +[Official appendix](https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html). + +Consulted: architecture-dependent feature and technical-limit tables. Supports +target-specific admission rather than universal device assumptions. +Individually verified footer: 2026-09-10. Access: 2026-10-21. No hardware +support matrix for Rheon was tested; FP64 throughput is not measured here. + +### GPU-C06. §2.7 NVCC: The NVIDIA CUDA Compiler + +[Official chapter](https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/nvcc.html). + +Consulted: §2.7.1 CUDA Source Files and Headers; §2.7.2 NVCC Compilation +Workflow. Supports separate host/device compilation, PTX, Cubin and fatbinary +artifacts. Accessed moving guide chapter on 2026-10-21; no individually checked +footer date is asserted for this page. Rust FFI choices in Chapter 38 are +proposed project decisions, not an NVIDIA-prescribed Rust integration. + +### GPU-C07. §5.5 Floating-Point Computation + +[Official appendix](https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/mathematical-functions.html). + +Consulted: §5.5.1 Floating-Point Introduction, normal/subnormal and rounding +discussion, and mathematical-function behavior. Supports declaring arithmetic +policy separately from storage types and checking numerical differences. +Individually verified footer: 2026-09-10. Access: 2026-10-21. The URL's filename +differs from the visible chapter title; the visible title above is intentional. + +### GPU-C08. §4.21 CUDA Interoperability with APIs + +[Official chapter](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/graphics-interop.html). + +Consulted: §4.21.2 External resource interoperability; §4.21.2.1 Vulkan +interoperability, including device matching, imported memory and +§4.21.2.1.9 Signaling/Waiting on Imported Synchronization Objects. Supports only +the optional render-interop discussion. Accessed moving guide chapter on +2026-10-21; no individually checked footer date asserted here. Interop is not +a dependency of either numerical backend. + +### GPU-C09. §1.3 The CUDA platform + +[Official chapter](https://docs.nvidia.com/cuda/cuda-programming-guide/01-introduction/cuda-platform.html). + +Consulted: §1.3.1 NVIDIA GPU compute capability; §1.3.2 CUDA Toolkit and NVIDIA +Driver; §1.3.3 PTX and §1.3.4 Cubins/Fatbins. Supports the CUDA platform/device +scope and distinction between driver, toolkit and image compatibility. Footer +individually verified 2026-09-10; accessed 2026-10-21. No GPU was enumerated, +driver installed or toolkit compatibility tested here. + +## Khronos Vulkan + +### GPU-V00. Documentation build identity + +[Official Vulkan documentation index](https://docs.vulkan.org/spec/latest/index.html), +“About,” build identity quoted by value in the revision policy above. + +Supports the exact consulted specification revision and generation date. +Access: 2026-10-21. “latest” is a moving URL. Specification 1.4.365 does not +imply a requirement for a Vulkan 1.4 device. + +### GPU-V01. Compute pipelines and dispatch + +[Shaders](https://docs.vulkan.org/spec/latest/chapters/shaders.html), introduction +and compute shader execution; [Dispatching Commands](https://docs.vulkan.org/spec/latest/chapters/dispatch.html), +`vkCmdDispatch`. + +Specification build: GPU-V00. Access: 2026-10-21. Supports compute workgroup +execution and dispatch semantics. Chapter 38's shader compilation and numerical +pipeline are proposed; no SPIR-V modules were compiled or executed. + +### GPU-V02. Devices and Queues + +[Official specification chapter](https://docs.vulkan.org/spec/latest/chapters/devsandqueues.html), +“Physical Devices,” physical-device properties and queue-family discovery. + +Specification build: GPU-V00. Access: 2026-10-21. Supports independent Vulkan +device enumeration and compute queue selection. Does not establish which +devices or queue families exist in the execution environment. + +### GPU-V03. Descriptor Sets + +[Official specification chapter](https://docs.vulkan.org/spec/latest/chapters/descriptorsets.html), +“Descriptor Set Layout,” descriptor bindings and descriptor-set use/update +requirements. + +Specification build: GPU-V00. Access: 2026-10-21. Supports explicit shader +resource bindings and lifetime-conscious descriptor use. Storage buffer layout +and per-attempt descriptors in Chapter 38 are proposed design choices. + +### GPU-V04. Features: query, enable, and shaderFloat64 + +[Official specification chapter](https://docs.vulkan.org/spec/latest/chapters/features.html), +opening feature-query/enable rules; `VkPhysicalDeviceFeatures`, +`shaderFloat64`; subgroup extended-type and synchronization2 feature entries. + +Specification build: GPU-V00. Access: 2026-10-21. Supports explicit +`vkGetPhysicalDeviceFeatures2` querying and logical-device enablement, and +optional f64 shader arithmetic. API version does not itself establish that an +optional arithmetic/atomic feature is enabled. + +### GPU-V05. Limits and subgroup properties + +[Official specification chapter](https://docs.vulkan.org/spec/latest/chapters/limits.html), +`VkPhysicalDeviceLimits`, `VkPhysicalDeviceSubgroupProperties`, +`VkSubgroupFeatureFlagBits` and subgroup-size-control properties. + +Specification build: GPU-V00. Access: 2026-10-21. Supports workgroup, storage +buffer and allocation-limit admission, subgroup size/stages/operations queries, +and avoidance of fixed-width subgroup assumptions. Extended operand-type +requirements also need GPU-V04; arithmetic-bit support alone is insufficient. + +### GPU-V06. Memory Allocation and mapped-memory visibility + +[Official specification chapter](https://docs.vulkan.org/spec/latest/chapters/memory.html), +device memory, allocation/memory types, mapped memory, +`vkFlushMappedMemoryRanges` and `vkInvalidateMappedMemoryRanges`. + +Specification build: GPU-V00. Access: 2026-10-21. Supports compatible memory +selection, explicit host coherence handling and the need for device completion +before invalidation/readback. Mapping alone does not invalidate noncoherent +memory. Chapter 38's byte budget and staging rings are engineering proposals. + +### GPU-V07. Floating-point atomics are separate features + +[VkPhysicalDeviceShaderAtomicFloatFeaturesEXT](https://docs.vulkan.org/refpages/latest/refpages/source/VkPhysicalDeviceShaderAtomicFloatFeaturesEXT.html), +“Description,” and the matching feature definitions in +[Features](https://docs.vulkan.org/spec/latest/chapters/features.html). + +Access: 2026-10-21. Standalone reference-page version is not separately printed; +specification chapter revision is GPU-V00. Supports separate buffer/shared +f32/f64 atomic and atomic-add feature bits, query and enable rules. Baseline +hierarchical reductions avoid requiring these optional additions. + +### GPU-V08. Shader scopes and arithmetic group operations + +[Shaders](https://docs.vulkan.org/spec/latest/chapters/shaders.html), +“Workgroup,” “Subgroup,” “Group Operations,” and “Arithmetic Group Operations.” + +Specification build: GPU-V00. Access: 2026-10-21. Supports scope limits and +implementation-dependent subgroup arithmetic evaluation order. A proposed +deterministic reduction must state its own tree and arithmetic policy; a +subgroup reduction is not a universal bitwise-reproducibility primitive. + +### GPU-V09. Synchronization and Cache Control + +[Official specification chapter](https://docs.vulkan.org/spec/latest/chapters/synchronization.html), +“Execution and Memory Dependencies,” “Pipeline Barriers,” fences, semaphores, +buffer barriers and queue-family ownership transfers. + +Specification build: GPU-V00. Access: 2026-10-21. Supports execution versus +availability/visibility distinctions, resource-hazard dependencies and the +legacy/synchronization2 design discussion. Chapter 38's dependency table is a +proposed resource schedule, not a validated command buffer. + +### GPU-V10. Floating-point controls + +[VkPhysicalDeviceFloatControlsProperties](https://docs.vulkan.org/refpages/latest/refpages/source/VkPhysicalDeviceFloatControlsProperties.html), +“Description,” including per-width signed-zero/Inf/NaN, denormal and rounding +properties. + +Access: 2026-10-21. Standalone reference-page version is not separately printed. +Supports querying admissible float-control behavior; availability of f64 +arithmetic is a distinct GPU-V04 requirement. No cross-backend bitwise +equivalence follows from the existence of these properties. + +## Repository evidence inspected without execution + +Main baseline: `fd7c0f0`. Existing project authorship is Puma's; this research +does not replace it. The +[CPU oracle inventory](../implementation/gpu-cpu-oracle-inventory.md) provides +the full source inventory and qualification limits. + +| Source | Observed constraint used in Chapter 38 | +| --- | --- | +| [`src/simulation.rs`](../../../src/simulation.rs) | f32 named-axis fields; candidate/accepted separation; host-slice `StateView`; transactional step semantics; actual stored-velocity divergence gate | +| [`src/pressure.rs`](../../../src/pressure.rs) | Jacobi/SGS PCG identities; f64 pressure/reduction arithmetic; true residual and physical qualification | +| [`src/solver_provider.rs`](../../../src/solver_provider.rs) | Synchronous host-slice pressure provider; provider produces a candidate that Rheon qualifies; bounded allocation and cancellation responsibilities | +| [`src/provider_step.rs`](../../../src/provider_step.rs) | Whole-step acceptance boundary and separate provider workspace accounting | +| [`src/recording/provider/runtime.rs`](../../../src/recording/provider/runtime.rs) | Accepted-context lineage; bounded runtime custody; explicit exclusion of restart, authentication and preemption guarantees | + +No hardware discovery, scientific workload, GPU workload, dependency install or +full build was performed for this bibliography. A direct shell HTTP retrieval +attempt was unavailable because the environment's network proxy returned 403; +official pages were successfully read through the available web retrieval tool. +No restriction was bypassed. The unresolved evidence is implementation and +hardware qualification, not source access. From db6f33d4f57b0033cf866059f65f7b2033837108 Mon Sep 17 00:00:00 2001 From: Jeremy Date: Wed, 21 Oct 2026 15:34:17 +0000 Subject: [PATCH 2/3] docs(research): pin GPU inventory links outside site bundle --- .../gpu-grid-publication/execution-ledger.md | 12 ++++ docs/plans/gpu-grid-publication/issues.md | 1 + docs/plans/gpu-grid-publication/plan.md | 6 +- .../reports/independent-publication-review.md | 56 +++++++++++++++++++ .../gpu-cpu-oracle-inventory.md | 10 ++-- 5 files changed, 77 insertions(+), 8 deletions(-) diff --git a/docs/plans/gpu-grid-publication/execution-ledger.md b/docs/plans/gpu-grid-publication/execution-ledger.md index 6eb227e5..bf52f524 100644 --- a/docs/plans/gpu-grid-publication/execution-ledger.md +++ b/docs/plans/gpu-grid-publication/execution-ledger.md @@ -42,3 +42,15 @@ check receipt. No production/proof/workflow/builder, FEM, portal or thermal file diff exists against the publication base. Generated checks remain outside Git. Real normal CI, generated PDF inspection and browser/source freshness are still required before parent integration and are not predeclared passed here. + +## Actual site gate correction + +Draft PR150 head `ab41f23` passed dense3D CI but compressible run +`39625662383` failed generated-site verification after its real Chromium PDF +render. Five inventory links targeted repository `experiments/` and `.github/` +files that the existing site source copier does not include. Relative file +existence checks therefore did not establish generated-site availability. +Changed those five authored links to immutable GitHub URLs at the original +`fd7c0f0` inventory epoch. No shared builder or workflow gate was changed. +The first failed run remains evidence; corrected-head CI and actual artifact +inspection are required. The successful render is not a visual review. diff --git a/docs/plans/gpu-grid-publication/issues.md b/docs/plans/gpu-grid-publication/issues.md index 51274240..1ba64998 100644 --- a/docs/plans/gpu-grid-publication/issues.md +++ b/docs/plans/gpu-grid-publication/issues.md @@ -9,3 +9,4 @@ | GPU-P05 | Full pipeline includes native/scientific qualification and actual PDF/browser gates | Newly authorized publication checks, normal workflows unchanged; no weaker substitute or GPU-science claim | | GPU-P06 | Older Library upload remains uncertain | Preserve all artifacts/identity state, no retry or alternate route; unrelated to repository publication | | GPU-P07 | Older executor may lack fresh artifact/site domain access | Exact target/status to parent fresh verifier if needed; no denied read retry or settings/security change | +| GPU-P08 | Actual site gate found five valid repository links outside the site copier scope | Pin those inventory links to immutable GitHub source URLs; retain failed run39625662383 and require corrected-head CI/artifact review | diff --git a/docs/plans/gpu-grid-publication/plan.md b/docs/plans/gpu-grid-publication/plan.md index ff0514bb..9623edb6 100644 --- a/docs/plans/gpu-grid-publication/plan.md +++ b/docs/plans/gpu-grid-publication/plan.md @@ -2,8 +2,8 @@ Status: Active documentation publication preparation. Current phase: normal exact-head publication CI and artifact-review preparation. -Next slice: push the reviewed publication head normally and create its draft PR -for exact-head CI and actual artifact review. +Next slice: review and push the five-link generated-site correction to draft +PR150, then obtain corrected-head CI and actual artifact review. ## Authority, source and ownership @@ -64,7 +64,7 @@ current-main reconciliation and explicit conflict reporting before integration. | --- | --- | --- | | P0 | Isolate from current main; preserve research/FEM/portal/product history; source/output/standards/build-scope review | Complete | | P1 | Prose-only GPU chapters, current CPU refresh, source/license/access audit and minimal discovery paragraph; independent source/math review | Complete; no deciding source findings remain | -| P2 | Exact immutable head pushed normally; draft PR; existing matching native/scientific/book/browser CI | Active preparation | +| P2 | Exact immutable head pushed normally; draft PR; existing matching native/scientific/book/browser CI | PR150 open; first site gate failed; corrected-head checks required | | P3 | Actual generated PDF visual inspection, GPU HTML/Markdown/source links, math, no code blocks, PDF/source freshness and browser evidence | Pending | | P4 | Parent handoff with exact PR/head/review/CI/artifact evidence and remaining blockers; no worker merge | Pending | diff --git a/docs/plans/gpu-grid-publication/reports/independent-publication-review.md b/docs/plans/gpu-grid-publication/reports/independent-publication-review.md index c28f0a48..d0850a65 100644 --- a/docs/plans/gpu-grid-publication/reports/independent-publication-review.md +++ b/docs/plans/gpu-grid-publication/reports/independent-publication-review.md @@ -159,3 +159,59 @@ installation, security change, full book build or deployment was run by this reviewer. The reviewer has not edited any chapter, bibliography, production source, shared navigation, build or workflow file. Any later normal pipeline or deployment evidence belongs to its actual runner and exact revision. + +## Generated-site link correction review + +Review date: 2026-10-21, after publication head +`ab41f23cc9fe3950965c028413898a0a23e021b0`. The publication ledger records that +real generated-site verification in run `39625662383` failed on five inventory +links after PDF rendering. This review inspects the corrective source diff; it +does not independently certify that CI run, a visual PDF review or the corrected +site output. The earlier source-only link check did not establish generated-site +availability, and its stated limitation remains material. + +**Correction disposition: no deciding blocker in the five-link repair.** The +inventory diff changes only five destination strings, keeping all labels, +prose, equations, whitespace and other content byte-identical after those exact +substitutions are reversed. Three experiment links and two workflow links now +use `https://github.com/MrScripty/Rheon/blob/` followed by the full original +inventory pin `fd7c0f0aee63bc8710df99f9bb01154e7619b974` and their original +repository paths. Each target independently exists as a Git blob at that exact +historical pin: + +- `experiments/aligned_stokes.rs` +- `experiments/submarine_stokes.rs` +- `experiments/plate_coupling.rs` +- `.github/workflows/research-book-pages.yml` +- `.github/workflows/research-book-reviewed-pages.yml` + +Direct inspection of `docs/education/build.py:99` confirms that external URLs +are left intact. Its source-copy branch at line 116 accepts repository-root +families `docs`, `evidence`, `proofs`, `src`, `examples` and `tests`; it omits +`experiments` and `.github`. The unchanged fallback at line 122 leaves their +former relative links unrewritten. The repair therefore addresses the actual +copy-rule mismatch without broadening the copier or weakening verification. +Historical source availability is checked locally at the Git pin; successful +public URL access and corrected generated-site behavior remain output-level +checks rather than claims from this inspection. + +Before this report append, the entire corrective diff comprised the inventory +link changes and three publication process files only: plan, ledger and issues. +Those process changes retain the failed run/head and require corrected-head CI +and actual artifact inspection; they do not declare rendering to be visual +review or the old source checks to be successful site qualification. No other +chapter, bibliography, production, workflow, builder, dependency, FEM, portal or +thermal file changed. Git whitespace checking passed. This reviewer appended +only this section and did not commit or push. + +The original nine-hash table above identifies the pre-correction reviewed +snapshot. The following corrected inventory identity supersedes that one row; +the other eight reviewed source identities remain unchanged: + +| Corrected path relative to repository | SHA256 | +| --- | --- | +| `docs/research-book/implementation/gpu-cpu-oracle-inventory.md` | `7c07f46261ae422fe2c5535d10b5e4718874be1ca9f500659b67955110af8687` | + +New committed-head identity confirmation and corrected same-head CI/site/PDF/ +browser evidence remain pending. No reviewer action was denied or left with +an uncertain mutation outcome. diff --git a/docs/research-book/implementation/gpu-cpu-oracle-inventory.md b/docs/research-book/implementation/gpu-cpu-oracle-inventory.md index e52b11f8..431b2a17 100644 --- a/docs/research-book/implementation/gpu-cpu-oracle-inventory.md +++ b/docs/research-book/implementation/gpu-cpu-oracle-inventory.md @@ -336,9 +336,9 @@ Column shear/momentum and the [flat MAC bridge](../../../src/column_mac.rs), mechanisms. Their topology, velocity families and parameter restrictions must be carried into any comparison. Adding an extruded third velocity component does not establish arbitrary three-dimensional geometry or flow. -The opt-in [aligned Stokes](../../../experiments/aligned_stokes.rs), -[submarine Stokes](../../../experiments/submarine_stokes.rs) and -[plate coupling](../../../experiments/plate_coupling.rs) experiments solve different +The opt-in [aligned Stokes](https://github.com/MrScripty/Rheon/blob/fd7c0f0aee63bc8710df99f9bb01154e7619b974/experiments/aligned_stokes.rs), +[submarine Stokes](https://github.com/MrScripty/Rheon/blob/fd7c0f0aee63bc8710df99f9bb01154e7619b974/experiments/submarine_stokes.rs) and +[plate coupling](https://github.com/MrScripty/Rheon/blob/fd7c0f0aee63bc8710df99f9bb01154e7619b974/experiments/plate_coupling.rs) experiments solve different restricted problems. They are not the default carrier or a general FEM backend. ## 7. Particle preview and display voxels @@ -486,11 +486,11 @@ Markdown/source downloads. [PDF freshness](../../education/pdf_freshness.py) binds source inputs, and [browser qualification](../../education/verify_browser.py) renders/qualifies the PDF and interactive behavior. -The [full research-book workflow](../../../.github/workflows/research-book-pages.yml) +The [full research-book workflow](https://github.com/MrScripty/Rheon/blob/fd7c0f0aee63bc8710df99f9bb01154e7619b974/.github/workflows/research-book-pages.yml) executes numerical references, native contracts/examples, scientific packet qualification, browser/Wasm work and render/publication checks; it is not a safe substitute for a source-only documentation check under this task's scope. -The [reviewed publication workflow](../../../.github/workflows/research-book-reviewed-pages.yml) +The [reviewed publication workflow](https://github.com/MrScripty/Rheon/blob/fd7c0f0aee63bc8710df99f9bb01154e7619b974/.github/workflows/research-book-reviewed-pages.yml) is another publication entry point and likewise is not evidence of a local run. No full build was attempted here. Inclusion of a note must not be described as successful PDF/Word/browser delivery until those actual outputs and freshness From 2da749bd08b7de44990b345b3b807736f18b8487 Mon Sep 17 00:00:00 2001 From: Jeremy Date: Wed, 21 Oct 2026 17:23:22 +0000 Subject: [PATCH 3/3] fix(ci): budget complete publication browser probes --- .github/workflows/research-book-pages.yml | 7 +- .../gpu-grid-publication/execution-ledger.md | 38 +++++ docs/plans/gpu-grid-publication/issues.md | 1 + docs/plans/gpu-grid-publication/plan.md | 38 ++++- .../reports/ci-budget-independent-review.md | 120 ++++++++++++++++ .../reports/ci-timeout-diagnosis.md | 135 ++++++++++++++++++ 6 files changed, 329 insertions(+), 10 deletions(-) create mode 100644 docs/plans/gpu-grid-publication/reports/ci-budget-independent-review.md create mode 100644 docs/plans/gpu-grid-publication/reports/ci-timeout-diagnosis.md diff --git a/.github/workflows/research-book-pages.yml b/.github/workflows/research-book-pages.yml index 1b31c0e9..e4854ae9 100644 --- a/.github/workflows/research-book-pages.yml +++ b/.github/workflows/research-book-pages.yml @@ -118,9 +118,10 @@ jobs: path: ${{ runner.temp }}/volume-edition/volume-evidence/ build: runs-on: ubuntu-24.04 - # Full Chromium publication probes exceeded 60m on run 38450395780. - # Keep every test; allow software-renderer variation and expose test names. - timeout-minutes: 90 + # Run 39626898860 was cancelled near 90m during late rejection probes. + # Repeated full Chromium traversals took 20–21m; retain every check within + # a 120m wall-clock budget. Per-case scientific limits stay unchanged. + timeout-minutes: 120 env: PYTHONDONTWRITEBYTECODE: '1' # Ubuntu's packaged AppArmor profile permits Chrome's sandbox at this path. diff --git a/docs/plans/gpu-grid-publication/execution-ledger.md b/docs/plans/gpu-grid-publication/execution-ledger.md index bf52f524..7069888b 100644 --- a/docs/plans/gpu-grid-publication/execution-ledger.md +++ b/docs/plans/gpu-grid-publication/execution-ledger.md @@ -54,3 +54,41 @@ Changed those five authored links to immutable GitHub URLs at the original `fd7c0f0` inventory epoch. No shared builder or workflow gate was changed. The first failed run remains evidence; corrected-head CI and actual artifact inspection are required. The successful render is not a visual review. + +## Preview acceptance and full-job timeout diagnosis + +Corrected PR150 headdb6f33d passed exact-head source review and preview run39626898869. +Actual artifact12357433282 binds its297-page PDF, SHA256 +`5bb728859fab3bb25dacc8591d52a0aff2f93168f26471dda84a9836179edc97`; +independent visual review covered everyGPU page161–192 and ten full pages. +Read-only artifact checks verified153 inputs against authored/build source or +retained generated figures, plus581GPU HTML asset/anchor targets and eight +packagedGPU Markdown documents. Local Chromium sandbox launch was blocked before +startup without a settings change or unsandboxed retry; parent verifier later +confirmed hosted behavior requirements. These results retain their own identities. + +Full run39626898860 did not pass: build119314799555 ran90m29s and was cancelled +during late runtime rejection probes. The actual render/site/retained browser +check passed; missing terminal regressions/aligned-strain suites and skipped +full-site artifact remain unsatisfied. Read-only permitted logs/metadata support +a wall-clock timeout strongly, with unavailable annotation noted. Three complete +browser traversals consistently took20–21minutes. A bounded publication-job +cap90→120 is proposed for parent shared-workflow coordination, preserving every +command, threshold, per-method limit and sandbox. No workflow mutation, rerun, +push or merge occurred during diagnosis. See the new detailed timing report. + +## Parent-coordinated bounded publication trial + +The parent explicitly authorized the independently reviewed90→120-minute +publication-job cap/comment and owned diagnosis/admission records, with every +command, numerical threshold, per-case limit, sandbox and test unchanged. +Applied exact candidate patch SHA256 +`7a950a9ee83d86d46e941f5ae16a46279e07670e72fb7591aa70298b13a0f58e`; +the resulting workflow SHA256 is +`b1f6b5afd3be0e725808ddd9857460d5ebe111c68ee87061aeb8464dbc214f48`. +Exact replacement normalization proves the workflow differs only in that +comment and one timeout scalar. The independent pre-application review is +retained unchanged in reports/ci-budget-independent-review.md. The old cancelled +run and missing annotation limit remain explicit. New exact-head normal CI and +a retained full-site artifact are pending; no manual rerun or merge is made. +Thermal PR149, GPU chapter/source hashes and all preserved branches are untouched. diff --git a/docs/plans/gpu-grid-publication/issues.md b/docs/plans/gpu-grid-publication/issues.md index 1ba64998..f8e69c4a 100644 --- a/docs/plans/gpu-grid-publication/issues.md +++ b/docs/plans/gpu-grid-publication/issues.md @@ -10,3 +10,4 @@ | GPU-P06 | Older Library upload remains uncertain | Preserve all artifacts/identity state, no retry or alternate route; unrelated to repository publication | | GPU-P07 | Older executor may lack fresh artifact/site domain access | Exact target/status to parent fresh verifier if needed; no denied read retry or settings/security change | | GPU-P08 | Actual site gate found five valid repository links outside the site copier scope | Pin those inventory links to immutable GitHub source URLs; retain failed run39625662383 and require corrected-head CI/artifact review | +| GPU-P09 | Full publication job cancelled at90m during late regression probes; full-site upload skipped | Terminal gate remains unsatisfied; measured timing and independent review support parent-authorized build-job cap120m trial; patch applied, new exact-head normal CI pending; no unchanged blind rerun | diff --git a/docs/plans/gpu-grid-publication/plan.md b/docs/plans/gpu-grid-publication/plan.md index 9623edb6..6f2d2bd7 100644 --- a/docs/plans/gpu-grid-publication/plan.md +++ b/docs/plans/gpu-grid-publication/plan.md @@ -1,9 +1,10 @@ # GPU research publication continuation Status: Active documentation publication preparation. -Current phase: normal exact-head publication CI and artifact-review preparation. -Next slice: review and push the five-link generated-site correction to draft -PR150, then obtain corrected-head CI and actual artifact review. +Current phase: authorized budget-only publication CI trial. +Next slice: commit and push the reviewed 120-minute build-job cap normally, +verify the exact head/tree and observe one matching normal CI run; retain its +terminal results and full-site artifact before any parent merge. ## Authority, source and ownership @@ -27,7 +28,10 @@ Write set: GPU Chapters 37–39; GPU implementation/source/license notes and bibliographies; original research plan history as inherited unchanged; this new publication plan/reports; and one additive GPU paragraph in the book README. No production, proof, Cargo, workflow, admission manifest or builder -edit is planned. FEM 30–36, portal 40 and thermal-owned paths stay unchanged. +edit was planned for the publication prose. The parent-authorized timeout continuation changes +only the shared workflow build-job budget/comment; +tests, ordering, sandbox and numerical limits remain unchanged. FEM 30–36, +portal 40 and thermal-owned paths stay unchanged. Any later shared README collision is reported to the parent's serial integration owner rather than overwritten. @@ -64,9 +68,9 @@ current-main reconciliation and explicit conflict reporting before integration. | --- | --- | --- | | P0 | Isolate from current main; preserve research/FEM/portal/product history; source/output/standards/build-scope review | Complete | | P1 | Prose-only GPU chapters, current CPU refresh, source/license/access audit and minimal discovery paragraph; independent source/math review | Complete; no deciding source findings remain | -| P2 | Exact immutable head pushed normally; draft PR; existing matching native/scientific/book/browser CI | PR150 open; first site gate failed; corrected-head checks required | -| P3 | Actual generated PDF visual inspection, GPU HTML/Markdown/source links, math, no code blocks, PDF/source freshness and browser evidence | Pending | -| P4 | Parent handoff with exact PR/head/review/CI/artifact evidence and remaining blockers; no worker merge | Pending | +| P2 | Exact immutable head pushed normally; draft PR; existing matching native/scientific/book/browser CI | PR150 draft atdb6f33d; preview/dense3D/lessons passed; full build cancelled at90m | +| P3 | Actual generated PDF visual inspection, GPU HTML/Markdown/source links, math, no code blocks, PDF/source freshness and browser evidence | Preview artifact12357433282 independently visually reviewed; full-site artifact and terminal regression evidence unavailable | +| P4 | Parent handoff with exact PR/head/review/CI/artifact evidence and remaining blockers; no worker merge | Active; timeout remedy proposed for shared-workflow coordination | Acceptance is publication-specific: source/math/license scope, original equation retention, no pseudocode, correct current-source limits, exact-head normal CI, @@ -99,5 +103,25 @@ uses a normal corrective/revert PR with new artifact identity and unchanged numerical gates, not mutation of retained outputs or history. See [ledger](execution-ledger.md), [issues](issues.md), +[budget-only review](reports/ci-budget-independent-review.md), [CPU refresh](reports/cpu-refresh.md) and [independent publication review](reports/independent-publication-review.md). + +## CI cancellation continuation + +The [timing diagnosis](reports/ci-timeout-diagnosis.md) records full run39626898860 +with build-job duration90m29s and cancellation during final regression probes. +Actual render/site/retained-check results passed, but the full gate remains +unsatisfied and no final full-site artifact exists. The parent explicitly coordinated and authorized the independently reviewed +120-minute wall-clock cap trial. The exact budget/comment patch is applied; +new-head CI remains pending. No manual duplicate rerun is planned. +It preserves all test commands, scientific work limits and sandbox behavior. +If one corrected run still exhausts the cap, stop and re-plan instead of +automatically increasing it again. Endpoint allowlist denials and cancellation +annotation uncertainty remain explicit in the diagnosis. + +The new trial is not acceptance of the cancelled run. GPU source review and +preview visual evidence keep their original db6f33d identities; new exact-head +CI and artifact evidence are required for this workflow revision. Thermal PR149 +and all branches remain untouched. Parent owns later normal integration and +merge order, with no merge until required gates pass. diff --git a/docs/plans/gpu-grid-publication/reports/ci-budget-independent-review.md b/docs/plans/gpu-grid-publication/reports/ci-budget-independent-review.md new file mode 100644 index 00000000..63ad4b67 --- /dev/null +++ b/docs/plans/gpu-grid-publication/reports/ci-budget-independent-review.md @@ -0,0 +1,120 @@ +# Independent review of publication CI cancellation diagnosis + +Reviewed 2026-10-21 against PR150 source head +`db6f33d4f57b0033cf866059f65f7b2033837108`. Read-only run: +`39626898860`; completed build job: `119314799555`. + +**Judgment: the diagnosis is supported, and one coordinated 90-to-120-minute +job-budget trial is a reasonable smallest remedy.** This is not a completed +gate, a guarantee of completion within 120 minutes, or permission for this +reviewer to apply/dispatch it. No workflow or CI architecture change was made. + +## Independent evidence checked + +The reviewer successfully retrieved the completed job's decoded logs through +`github_fetch_workflow_job_logs` and current job conclusions through +`github_fetch_workflow_run_jobs`, both supported read-only connector actions. +The latter reports build `cancelled`, bounded-lessons `success` and deploy +`skipped`. The full decoded job log contains 2,551 lines and corroborates the +recorded source checkout and all decisive subcommand milestone timestamps. +The saved `timing-evidence.json` supplies the job/step start/completion summary; +its elapsed-time arithmetic is consistent with those milestone logs. + +The configured cap is 90 minutes at +`.github/workflows/research-book-pages.yml:123`. The recorded job interval is +15:35:23–17:05:52 UTC, or 5,429 seconds including cleanup. The log emits +`The operation was canceled` at 17:05:49.0802842Z and immediately performs post-job +cleanup. No assertion failure or traceback appears before that cancellation. +All qualification steps before the final grouped browser-runtime step have +successful terminal conclusions; the full-site artifact upload was skipped. + +Independently recomputed intervals are: + +| Observed interval | Seconds | Interpretation | +| --- | ---: | --- | +| 15:44:58.994–16:06:21.756 | 1,282.762 | Build receipt to successful actual render-pdf result; includes browser qualification, not PDF printing alone | +| 16:10:43.502–16:31:57.081 | 1,273.579 | Explicit retained browser check | +| 16:31:57.081–16:52:10.645 | 1,213.564 | Runtime positive test including suite startup | +| 16:52:20.565–17:05:49.080 | 808.515 | Observed interrupted interval attributed to the next missing-hub test by ordering; no separately logged start timestamp | + +The HTTP and JavaScript rejection results are present at 16:52:15.2059708Z and +16:52:20.5654326Z. The final runtime suite and both subsequent aligned-strain +suite invocations have no successful terminal result in this run. Passing prior +render/site checks or the separate preview artifact does not close those gates. + +The official [GitHub timeout definition](https://docs.github.com/en/actions/reference/workflows-and-actions/workflow-syntax#jobsjob_idtimeout-minutes) +states that the job cap controls automatic cancellation. The approximately +90-minute elapsed time strongly supports that cause. The accessible log/API +does not establish the deciding timeout annotation, so manual or platform +cancellation remains possible. The diagnosis correctly retains this uncertainty. + +## Source control flow checked + +`docs/education/test_browser_runtime.py:32` invokes the same complete +`verify_browser.verify` entry point for all six probes and checks retained +PDF/receipt bytes afterward. A probe may still fail early; entry-point identity +does not mean every negative probe finishes every qualifier. + +`verify_browser.py:131–139` runs teaching, transport, timestep and flow before +native qualification. `native_browser.py:21–22` checks the mutated hub only +after loading `native-labs.html`. Thus the missing-hub test's long repeated +native-flow traversal is explained by actual order, rather than evidence of a +new GPU workload or a demonstrated transient stall. Default alphabetical test +ordering places it after the three completed tests and before mobile-error and +unresponsive-control probes. + +The flow qualifier keeps two pressure methods and 80 accepted steps per method; +`flow_browser.py:841` retains the 600-second deadline for each method's resumed +trajectory loop. This is a bound on that phase, not the entire browser process. +The remaining mobile-error/control mutations can be detected before the long +flow path (`verify_browser.py:68–71`, `test_browser_runtime.py:51–61`). The +aligned-strain browser suites invoke their direct strain qualifier +(`test_aligned_strain_browser.py:60–78`), not the full flow traversal. Their +remaining completion time was not measured by this cancelled run. + +The historical 60-to-90-minute cap change at commit +`900e30b5807497c3c2bb20e3d7d0fed6d9e42cc1` was inspected as source history. It +does not supply a comparable successful timing or prove the next cap sufficient. + +## Final unapplied candidate review + +Candidate: `budget-only-candidate.patch` in this scratch directory. +Final patch SHA256: +`7a950a9ee83d86d46e941f5ae16a46279e07670e72fb7591aa70298b13a0f58e`. +The candidate targets only `.github/workflows/research-book-pages.yml`. +Removing comment lines and reversing the single build timeout scalar from 120 +to 90 reproduces the existing workflow exactly. The candidate therefore retains +every command, order, test, assertion, numerical threshold, pressure method, +trajectory length, per-method deadline, sandbox setting, artifact/deployment +dependency and error policy. A read-only `git apply --check` passed. + +Original workflow SHA256: +`a4ab0fdddbd141d0323db8a2dd800c484c235dd23121028f5b9d14bfe20acc67`. +Prospective workflow SHA256 if that exact patch is applied: +`b1f6b5afd3be0e725808ddd9857460d5ebe111c68ee87061aeb8464dbc214f48`. + +One accuracy correction was suggested and accepted: the explanatory comment +now says the run was **cancelled near 90 minutes**, rather than stating an +unconfirmed timeout cause as certain. No deciding review finding remains open. +The extra 30 minutes is headroom for the observed unfinished tail plus remaining +work, not an empirically established upper bound. One changed-cap run can test +that decision without weakening acceptance. An unchanged rerun cannot test it; +parallelizing generated-input/native/browser stages would require a larger +same-head custody and dependency review. If the bounded trial still fails or +expires, stop and re-plan rather than raising the cap automatically. + +## Authorization and access boundary + +The shared-workflow owner must coordinate application before a new exact head +is committed/pushed. The resulting normal CI must retain all required terminal +results and full artifact evidence; no merge is justified by this diagnosis. +The shared workflow remains unchanged at this review point. This reviewer wrote +only this scratch report, did not modify the repository diagnosis/report or +workflow, and ran no CI, scientific, numerical, GPU or browser workload. + +The previously rejected check-run annotation and workflow-run-list endpoints +were not called, retried or reached through another route. Their rejection +leaves detailed cancellation annotation and historical-run comparison +unavailable. No reviewer action was denied or left with an uncertain mutation +outcome. Parent coordination, actual application and the changed-head CI result +are pending; the proposed patch is not an executed remedy. diff --git a/docs/plans/gpu-grid-publication/reports/ci-timeout-diagnosis.md b/docs/plans/gpu-grid-publication/reports/ci-timeout-diagnosis.md new file mode 100644 index 00000000..56a42be2 --- /dev/null +++ b/docs/plans/gpu-grid-publication/reports/ci-timeout-diagnosis.md @@ -0,0 +1,135 @@ +# Publication CI cancellation diagnosis and proposed bounded remedy + +Prepared 2026-10-21 for parent coordination. Initial diagnosis checkpoint: +shared-workflow remedy proposed, not applied or dispatched. PR150 remains draft +at `db6f33d4f57b0033cf866059f65f7b2033837108`. Its GPU book content and independent +preview-PDF review are unchanged. Cancellation is not a successful full gate. + +## Recorded cause and uncertainty + +Run [39626898860](https://github.com/MrScripty/Rheon/actions/runs/39626898860), +build job [119314799555](https://github.com/MrScripty/Rheon/actions/runs/39626898860/job/119314799555), +started `15:35:23Z` and completed `17:05:52Z`: 90 minutes 29 seconds including +cleanup. Its workflow sets `jobs.build.timeout-minutes: 90`. The last command +was cancelled at `17:05:49.080Z` with `The operation was canceled`; cleanup +followed immediately. The job conclusion is `cancelled`, not `failure` or +`success`. All earlier steps have successful terminal conclusions. No assertion +failure appears before cancellation. + +This is strong evidence for the configured job wall-clock timeout. GitHub's +[official timeout documentation](https://docs.github.com/en/actions/reference/workflows-and-actions/workflow-syntax#jobsjob_idtimeout-minutes) +states that the job is automatically cancelled at its configured limit. The +accessible run/job metadata and log do not contain the deciding timeout +annotation, so manual/platform cancellation cannot be excluded absolutely. +The connector rejected the check-run annotation endpoint as an unsupported URL; +no alternate route was tried. A workflow-run-list endpoint was also rejected; +no alternate listing route or historical timing claim was substituted. + +## Measured per-step and subcommand timing + +All times below are runner log/API timestamps on 2026-10-21, not locally +executed workloads or a performance claim about the solver. + +| Phase | UTC interval | Elapsed | Terminal evidence | +| --- | --- | --- | --- | +| Setup through native browser fixtures | 15:35:23–15:43:48 | 8m25s | Successful | +| Static edition, PDF and site/lesson verification | 15:43:48–16:10:43 | 26m55s | Successful | +| Build receipt emitted to actual PDF render result | 15:44:58.994–16:06:21.756 | 21m22.8s | Render result, no page/network errors | +| Site verification | render result–16:06:33.766 | about12s | 47sections,1776math,9010localtargets;918ZIP links;467portablePDFlinks | +| Additional lesson probes and optimized timestep suite | 16:06:33.766–16:10:43.467 | 4m09.7s | All passed | +| Explicit retained browser check | 16:10:43.502–16:31:57.081 | 21m13.6s | Passed, Chromium155.0.8059.84 | +| Runtime positive probe, including suite startup | 16:31:57.081–16:52:10.645 | 20m13.6s | `test_current_six_labs_pass_without_pdf_or_receipt_writes ... ok` | +| HTTP rejection probe | previous result–16:52:15.206 | 4.6s | Passed | +| JavaScript rejection probe | previous result–16:52:20.565 | 5.4s | Passed | +| Pending missing-native-hub rejection probe | previous result–17:05:49.080 | 13m28.5s observed | Interrupted; no terminal result | + +The rendered/site-verified book and explicit retained check therefore passed. +The final runtime suite did not finish. The other runtime probes and both +normal/optimized aligned-strain suites have no terminal evidence in this run. +The final full-site artifact upload was skipped. Preview artifact12357433282 +remains separate valid evidence for its own exact source and PDF. + +## Source explanation and alternatives + +`test_browser_runtime.py` calls the complete `verify_browser.verify` for every +probe, copying the edition and checking that retained PDF/receipts do not change. +The complete qualifier includes `flow_browser.qualify`, whose admitted native +Wasm trajectory exercises 80 accepted steps for each of two pressure methods. +This is existing browser qualification, not a GPU backend experiment. The +missing-native-hub mutation is checked by `qualify_native` only after teaching, +transport, timestep and flow qualification. Thus its negative test legitimately +repeats much of the roughly20-minute positive path before observing the error. +Alphabetic unittest order places that missing-hub test after the three reported +successful tests. Its start is inferred from that order and the last successful +result; no separately emitted start timestamp is claimed. + +The three completed full-browser traversals are consistently about20–21minutes; +there is no measured transient stall supporting an unchanged blind rerun. At +the observed rate, finishing the interrupted late negative path alone could +need roughly another7minutes, before remaining runtime and aligned-strain tests. +This is an estimate, not an observed completion or a guaranteed bound. +The repository already raised this same job from60 to90minutes in commit +`900e30b5807497c3c2bb20e3d7d0fed6d9e42cc1` after an earlier publication overrun; +that history establishes intent to retain coverage, not a comparable timing +measurement for this source. + +## Chosen smallest remedy for parent coordination + +Propose changing only this publication job's wall-clock cap from90 to120minutes, +with a comment naming this run and the measured repeated browser traversals. +The extra30minutes covers the observed interrupted tail plus bounded uncertainty +for the unmeasured remaining suites. It is a capped resource decision, not a +promise that120minutes must succeed. No new solver cohort, trajectory length, +iteration limit, acceptance threshold, per-method600-second bound, sandbox +setting, test, command order or deployment gate changes. No `continue-on-error`, +reduced fixture, numerical fallback or unqualified artifact acceptance is added. + +A parallel-job design would add generated-input and native-packet transfer, +same-head provenance, freshness restoration and final deployment dependencies. +Concurrent browsers on one runner could instead introduce CPU/memory contention. +Those are larger changes than this evidence currently justifies. They remain +re-plan options if one corrected bounded run still exhausts120minutes. + +Exact prospective write set: `.github/workflows/research-book-pages.yml` budget +and explanatory comment, plus this GPU publication plan/ledger/issues/report. +The parent must coordinate ownership of that shared workflow before application. +The full test commands and coverage remain byte-identical. Independent review +must check that normalized diff before a normal commit/push. A rerun at the old +head cannot test the changed cap because it uses the old workflow revision. + +After coordination: apply the reviewed budget-only change, record a new exact +head, let matching normal CI execute once and retain the terminal results and +full-site artifact. Do not manually dispatch unrelated successful cohorts or +retry this cancelled run unchanged. Preserve prior failures and preview/source +reviews with their original identities. If the120-minute run fails, cancels or +times out, stop and re-plan rather than automatically raising the budget again. +No merge until every required gate completes and parent artifact/publication +verification is satisfied. This report executes no science or browser workload. + +## Access and execution status + +Read-only completed job logs, run metadata and job/step timestamps were retrieved +successfully through the permitted GitHub connector. The following requests +were rejected by the connector allowlist with HTTP400 and were not retried: +`/repos/MrScripty/Rheon/check-runs/119314799555` and +`/repos/MrScripty/Rheon/actions/workflows/research-book-pages.yml/runs?per_page=15&status=completed`. +Those rejected reads caused no mutation and leave the detailed cancellation +annotation and recent-run comparison unavailable. No workflow/source change, +CI rerun, push, merge, toolchain installation, local science, GPU workload, +browser launch or security change was performed for this diagnosis. + +## Authorized application checkpoint + +After the diagnosis and independent review, the parent explicitly authorized +the exact90-to-120-minute publication-job cap trial and explanatory comment, +plus these owned admission records. The reviewed candidate patch SHA256 is +`7a950a9ee83d86d46e941f5ae16a46279e07670e72fb7591aa70298b13a0f58e`. +Its applied workflow SHA256 is +`b1f6b5afd3be0e725808ddd9857460d5ebe111c68ee87061aeb8464dbc214f48`. +Normalization againstdb6f33d proves all other workflow bytes unchanged. +The [independent candidate review](ci-budget-independent-review.md) retains its +pre-application identity and access limits. A normal new-head commit/push and +one matching automatic CI run are the next execution, not a manual duplicate +rerun. Their actual head/tree/run identities belong in the handoff rather than +being predeclared here. No terminal success or new artifact is claimed at this +checkpoint. If this trial times out, stop and re-plan from its timings.