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 new file mode 100644 index 00000000..7069888b --- /dev/null +++ b/docs/plans/gpu-grid-publication/execution-ledger.md @@ -0,0 +1,94 @@ +# GPU publication execution ledger + +## 2026-10-21 — publication admission and source work + +Fetched main `d67aec1` after portal PR148 landed. Created a separate publication +branch/worktree and prepared a normal merge of preserved research `295e9be` with +no conflicts. Existing FEM, portal and product files remain unchanged. Read +current book build/output/publication instructions, FEM publication policy and +Coding-Standards `0fc92a2`; no ancestor/Rheon AGENTS or application skills were +present. Reused the existing CPU and independent reviewers; no duplicate worker +was created. + +Converted all nine GPU code/diagram blocks into prose/equations. Independent +Pandoc parsing confirmed every original 25 TeX expression retained, with 42 +current expressions and zero chapter CodeBlocks at that checkpoint. A dated CPU +refresh documents the unchanged pressure/operator baseline and actual new +authored-force/recording source contracts. New public license/source audit +distinguishes original exposition, scientific/API provenance and exact BSD +implementation comparisons without code/dependency import. One additive book +README paragraph supplies discovery; the automatic chapter/guide collector is +unchanged. Publication CI and actual artifact review remain pending. + +No GPU implementation/workload, hardware query, benchmark or production change +was made. Earlier Library artifacts and uncertain upload remain intact and were +not retried. Exact commit/PR/run/artifact evidence follows outside this +self-referential source snapshot. + +## Focused publication-source gate + +The reused independent reviewer found no unresolved deciding source/math/license +issue in the final prose revision. It retained all 25 original chapter TeX +expressions, with 42 current expressions, zero CodeBlocks and 181 resolving +relative file links. Root separately checked file/Markdown-anchor targets and +strict KaTeX 0.16.22 rendering for all 42 expressions, with no errors. The pinned +renderer assets were installed through the existing npm lock, scripts disabled, +into a task-owned external directory/cache; no project dependency was changed. +Existing `test_pdf_freshness.py` (10) and `test_publication_gates.py` (18) passed; +the latter uses browser fixtures and is not an executed Chromium check. + +Exact reviewed source hashes are in the independent report and external focused +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. + +## 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 new file mode 100644 index 00000000..f8e69c4a --- /dev/null +++ b/docs/plans/gpu-grid-publication/issues.md @@ -0,0 +1,13 @@ +# GPU publication issues and disposition + +| ID | Issue | Disposition | +| --- | --- | --- | +| GPU-P01 | Original research contains nine code/diagram blocks inconsistent with later no-pseudocode publication rule | Replaced with prose and recurrences; original math retained; independent final review required | +| GPU-P02 | Original `fd7c0f0` inventory predates authored-force/recording additions | Historical pin retained; dated `d67aec1` refresh integrated; no numerical result inferred | +| GPU-P03 | Scientific/API citation is distinct from reusable implementation source | Public scope/license audit and exact BSD comparison pins; no external code/figure/dependency imported; access limits retained | +| GPU-P04 | Concurrent thermal publication may touch shared book discovery | Parent serial integration; only one additive GPU README paragraph, preserve thermal and portal/FEM blocks; report conflicts | +| 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 | +| 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 new file mode 100644 index 00000000..6f2d2bd7 --- /dev/null +++ b/docs/plans/gpu-grid-publication/plan.md @@ -0,0 +1,127 @@ +# GPU research publication continuation + +Status: Active documentation publication preparation. +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 + +The user's “Get that published” authorizes normal publication preparation and +checks for the GPU research. Production GPU implementation, GPU/hardware +validation, benchmark/speedup claims and physics changes remain excluded. +The parent owns merge order and final lead merge/deployment decision; this worker +must report PR/head/review/CI/PDF/browser evidence first and does not merge. + +The preserved research branch is `research/gpu-grid-algorithms-20261021` at +`295e9be5736a92425b4281047576d959bbc1d5b2`. Publication branch +`documentation/gpu-research-publication-20261021`, worktree +`/workspace/rheon-gpu-publication`, starts at main +`d67aec1a4575e0a61aa8cd931c8eeec33c8dc6a3`, tree +`8e630e9f870ea3f59ac7962c557089d50bc6cfaf`. A normal merge preserves the research +commit in ancestry. Main includes FEM PR147, portal PR148 and product PR145/146. +No shared-index conflict occurred at this merge preparation. No rebase, force +push, deletion or rewriting of another branch is authorized. + +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 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. + +## Publication decisions and contract + +Puma remains the credited author. Replace nine old pseudocode/diagram blocks +with original explanatory prose and explicit mathematical recurrences, retaining +all original equations and numerical/termination conditions. Publish a public +source/license/access audit; use exact inspected BSD implementation comparison +pins, distinguish them from papers/proprietary API documentation used as factual +provenance, and incorporate no upstream code/dependency/figure or substantial +quoted expression. Future reuse requires its own complete component audit. + +The original CPU pin remains `fd7c0f0`; a dated source refresh at publication main +documents new authored-field and force-recording custody without relabeling old +reviews or numerical receipts. Pressure/operator/geometry are unchanged in the +delta. Published GPU capability, device support and speed remain prospective. + +Composed-design review: applicable to the candidate independent-backend design, +whose complete ownership/change/dependency/replacement probe remains in the +preserved research plan and Chapter 39. This publication adds no runtime owners +or architecture. Its produced artifact is prose/equation HTML, reading PDF and +portable Markdown/source bundle, all from the existing book pipeline. No generic +registry, renderer, archive or native GPU adapter is introduced. + +Concurrent-plan-integration profile: not applicable to these separate single-writer +plans with a designated parent integration owner, rather than competing proposals +to the same mutable plan revision. Shared publication source still requires +current-main reconciliation and explicit conflict reporting before integration. + +## Milestones and acceptance claims + +| Milestone | Owned output and deciding gate | State | +| --- | --- | --- | +| 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 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, +actual PDF and browser/link/source fidelity. Status: pending. Source parse is +not a substitute for rendered output; green CI is not a GPU numerical result. +All generated plots, fields, logs, site, PDF, screenshots and archives remain +outside Git under repository policy. Remote jobs' own caps are 15/90/35 minutes +for book lessons/full edition/compressible preview; unfiltered dense3D also runs +with no explicit job timeout. Do not weaken or bypass their gates. + +The older Library PDF/Markdown and uncertain upload outcome remain untouched. +Repository publication is not a Library retry. If this executor cannot retrieve +artifacts or reach the publication site under effective access, report the exact +target/action/reason/outcome for the parent's existing fresh verifier; do not +retry old denied reads, alter sandbox/security or substitute stale output. + +Re-plan for changed pressure/force/recording semantics, an incompatible/unread +source claim, lost original math, output clipping, source/receipt mismatch, +shared-index overlap, failed CI or denied required verification. A failed check +remains failed/pending until the named evidence resolves it. + +## Retained resources and rollback + +The publication branch/worktree are retained for the draft PR and parent review. +After remote publication they are shared and are not rewritten. Parent owns +normal merge or rejection and eventual clean-resource retirement; no cleanup is +performed here. The original research ref and uncertain Library artifacts are +independently preserved. If published documentation needs correction, parent +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. diff --git a/docs/plans/gpu-grid-publication/reports/cpu-refresh.md b/docs/plans/gpu-grid-publication/reports/cpu-refresh.md new file mode 100644 index 00000000..0a3462e6 --- /dev/null +++ b/docs/plans/gpu-grid-publication/reports/cpu-refresh.md @@ -0,0 +1,272 @@ +# CPU oracle refresh for GPU book publication + +Source-only review, 21 October 2026. Original GPU research baseline: +`fd7c0f0aee63bc8710df99f9bb01154e7619b974`. Publication-base main: +`d67aec1a4575e0a61aa8cd931c8eeec33c8dc6a3`. This report inspected the actual +Git/source delta and proposes documentation corrections. It does not claim a +new numerical, runtime, device, benchmark or recording test result. No production +source, existing GPU document, science/build command or unrelated worker was +changed by this review. Root owns integration of the proposed text. + +## Change provenance and unchanged oracle + +PR145 merged at `7516a8e5d328c6dc8eda8ea577cb03f20b2e343a`. Its relevant source +commits are `f55cf595f8511fca9afc86cc572f26b017fb79ca` (validated authoring/evaluator), +`19e0145ce652fac565cd04b4e07ace3bbaa46954` (MAC/particle stage adapters), and +`8d2b4f8a6868b357df9e9be178c8fa290b7eda77` (accepted-force recording custody). +PR146 merged at `b9218038d138d1c73818cf4158757521853a5cd7`; it integrates native +desktop authoring, immutable request/run intent, dispatch and recording. The +selected core force/simulation/recording files have no additional delta between +those two merge results. PR147's separately owned FEM chapters and PR148's fluid +portal research are present at the publication base but confer no GPU or general +fluid implementation capability. + +The inspected delta leaves `geometry.rs`, `sampling.rs`, `advection.rs`, +`operator.rs`, `pressure.rs`, `sgs.rs`, `viscosity.rs`, `liquid_step.rs`, +`solver_provider.rs`, `session.rs` and `material_session.rs` unchanged. Thus the +original filled-box MAC layout, six-f64-array pressure workspace, integrated +operator sign/units, zero pressure initialization, Jacobi/SGS PCG identities, +full true-residual plus predicted-divergence conjunction, stored-f32 divergence +and Courant acceptance remain the first pressure oracle. This is source identity +evidence, not a rerun of those contracts. The initial dense sealed-box GPU +recommendation remains appropriate. + +The meaningful refresh is additional authored external-force inputs and their +accepted publication custody. The original inventory must retain its true +historical pin; a new dated refresh must distinguish publication-base source +observations from the original research and from historical execution receipts. + +## Actual authored-field contract + +[force_field.rs](../../../../src/force_field.rs) now owns version-1 `ForceScene` +authoring and validated `ForceSnapshot` evaluation. `freeze()` validates the +complete scene, including disabled objects, then clones it into an owned `Arc`. +It permits at most 256 objects with unique nonzero IDs. Objects carry name, +enabled state, world translation, unit quaternion, units, kind, support and time. +There is no scale transform or mutable evaluator access. Deserializing a scene +does not bypass the freeze boundary. + +The admitted kinds are directional local vectors, signed radial constant +magnitude (positive outward, zero at the center), and signed vortex constant +tangential magnitude around local +Y (zero on the axis). Vortex positive +handedness sends +X toward -Z. These are explicit normalized fields, not inverse +square forces or a vorticity-confinement stencil. Supports are global, local +sphere with open outer boundary, or local half-open box +`[-half_extent,+half_extent)`. An optional edge fraction supplies the existing +smoothstep weight. Evaluation transforms the query to local space, evaluates +support and field, then rotates the weighted result to world space. + +Activation is enabled and `start <= accepted_interval_start < end` (or no end). +Intervals are not split at activation endpoints. Consequently sampling at a +midpoint/end, integrating the on/off interval exactly, or advancing the field +time during particle stability retries would change the current CPU method. +Quaternion validation tolerates squared-norm error at most `1e-12`; evaluation +does not silently normalize that quaternion. Stable direction normalization +handles zero, tiny and large coordinates. The frozen object order defines +superposition order; a parallel object reduction would introduce a new rounding +order that needs an explicit policy. + +`ForceUnits::Acceleration` is m/s². `ForceDensity` is N/m³ and is divided by the +positive density supplied by the numerical owner, never by particle mass. The +common authoring/evaluation definition therefore does not make the two solver +adapters' density or sampling rules interchangeable. + +## MAC stage, timestep and stamp changes + +[simulation.rs:420](../../../../src/simulation.rs#L420) adds +`step_with_force_fields(requested_dt, source, fields, cancel)` and +`step_with_emitter_and_force_fields(requested_dt, emitter, fields, cancel)`. +Their internal construction explicitly supplies no tracer barrier, prescribed +boundary, free surface, viscosity or provider, and no legacy `BodyForce` slice. +They are closed-box carrier entry points, not arbitrary combinations of every +optional `Simulation` capability. The emitter remains its existing separate +source stage. + +[forces.rs:128](../../../../src/forces.rs#L128) evaluates the immutable scene at +each velocity component's own world-space MAC face position, using constant +`SimulationConfig.density` and accepted interval start. The stage follows +advection and legacy smoke/nozzle forcing and precedes pressure. Normal outer +faces remain unchanged. Writes use existing f32 storage and the force work +report uses the actual stored velocity change with the existing discrete face +mass. `BeforeForces` and per-z-slice `ForceSlice` cancellation remain cooperative. + +[simulation.rs:793](../../../../src/simulation.rs#L793) adds an authored-field +acceleration rate to timestep preflight. The snapshot bound sums active-object +magnitude bounds by triangle inequality, ignores cancellation/support overlap, +and includes a small rotation/rounding allowance. It is multiplied by +`sum(1/h_d)` before entering the existing acceleration displacement budget. A +GPU maximum sampled only on occupied faces is not the same preflight policy. +No new pressure equation or pressure tolerance results from this force addition. + +The new `AcceptedForceIntervalPressure` exposes pressure, actual interval start, +step report, source/emitter controls, frozen scene and optional external-stage +work. On authored success the legacy `accepted_interval_pressure()` stays absent +to avoid presenting incomplete controls. Both legacy and force-aware carrier +stamps are cleared before a later prepare attempt, reset or commit; even a +rejected next attempt invalidates pressure inspection while accepted velocity, +tracer, time and generation remain intact. An empty authored scene skips field +work in the MAC numerical path but still receives force-aware custody: explicit +empty metadata is distinct from absence of authored metadata. + +## Particle adapter and different rejected-attempt semantics + +[particle_fluid.rs:403](../../../../src/particle_fluid.rs#L403) adds +`ParticleFluidSimulation::step_with_force_fields`. After candidate emission and +the existing density/EOS evaluation, it adds scene acceleration to gravity, +then accumulates existing pair forces. Query positions are pre-integration +particle positions, including newly emitted particles. Force density divides by +the evaluated local SPH density used by that pressure stage, not the material's +reference density, constant MAC density or parcel mass. The field time remains +the accepted interval start on all up-to-sixteen stability retries. The existing +post-force acceleration/damping screen includes the resulting total force; +there is no new MAC-style field-bound preflight in this particle adapter. + +Committed particle diagnostics still refer to post-emission/pre-integration +evaluation, now including gravity, authored fields and pair forces. They exclude +subsequent wall/contact response. Accepted positions/velocities remain at the +new publication time. Pair kernels, EOS, neighbor visibility and contact geometry +are not replaced by authored forces and gain no continuum-water qualification. + +`AcceptedParticleForceControls` retains actual interval start, original requested +dt, emitter, actual accepted report and frozen scene only after success. Unlike +carrier pressure scratch, its diagnostic storage has independent ownership: +rejected/cancelled later attempts preserve the prior particle force publication. +A successful legacy particle step clears the authored stamp; a successful +authored step replaces it. Documentation must not generalize carrier pressure +stamp invalidation into a rule that all accepted particle metadata disappears +on any next attempt. The synchronous `MaterialSession` wrapper was not extended +to expose these new authored APIs in this delta. + +## Recording additions and preserved limits + +[recording/force_fields.rs](../../../../src/recording/force_fields.rs) adds a +versioned binary input document within the existing SQLite envelope. The SQL +schema version and legacy field encodings/hash algorithms are unchanged. +The new field is `external_force_scene_v1`, scalar u8, one component, dimensions +`[encoded_byte_count,1,1]`, units `1`, association +`immutable_run_input; world_m; activation_at_interval_start_s`, stage +`external_force_document_v1`, with no evaluation timestamp. It represents +immutable input metadata, not sampled acceleration or an endpoint force field. +Encoding is little-endian, versioned, capped at 256 KiB and preserves f64 bits, +including signed zero, full IDs/order, disabled objects and activation controls. +Decoding bounds bytes/count/name length, rejects invalid/trailing data and +freezes the reconstructed scene. + +`RecordingRun::with_force_fields` binds the complete document's SHA256 into the +existing configuration header. `validate_run_force_fields` requires either +legacy absence on both sides or matching run/step force metadata. Append and +selected read both apply it. Therefore the current force-aware archive binds +one immutable scene for the run: changing even a disabled object's name or +bits cannot be silently treated as the same run input. Absence means legacy +metadata, not proof that an arbitrary producer applied no force. This binding +checks integrity/association, not authenticated physical execution. + +[carrier_publication](../../../../src/recording/carrier.rs#L46) now accepts +either the legacy success stamp or the actual force-aware pressure stamp. The +legacy canonical inventory remains four constructor fields/five accepted fields; +an authored publication adds one metadata field. The adapter automatically +retains the actual accepted scene plus external-stage work/control provenance. +`carrier_publication_with_force_fields` attaches constructor metadata or checks +that an accepted publication already owns exactly the supplied scene. It cannot +relabel an unstamped accepted publication or substitute later draft force data. +The existing error string mentioning only `step/step_with_emitter` is narrower +than the new control flow; source branches and tests, not that string alone, +establish force-aware admission. + +[material_publication](../../../../src/recording/material.rs#L90) similarly +checks supplied requested-dt bits, emitter and report against actual accepted +particle-force custody, then retains the actual scene. Its explicit force-aware +wrapper attaches constructor metadata or requires actual accepted custody. +Both decoded carrier and material values now expose optional force metadata. +No derived per-particle acceleration is recomputed from the scene during readback. + +This is a new typed force-document codec, so the broad statement that *all* +recorded controls are opaque Debug text needs qualification. Most configuration, +report and source-control text remains opaque; the force document itself is now +bit-preserving validated typed input. It still is not a complete restart codec: +mesh bytes, source/emission/integrator state and all other required state remain +absent. SQLite acceptance-after-physics, durable-prefix/unrecorded-tail, +uncertain-append poisoning, keep-all limits and cold-selection semantics are +unchanged. Legacy `step_with_forces`, viscosity and arbitrary GPU producers do +not gain a recording stamp merely because authored-field adapters exist. + +PR146's [desktop dispatch](../../../../src/bin/desktop/carrier_execution.rs#L131) +uses these authored APIs for the native carrier and the direct particle owner. +Its [force eligibility policy](../../../../src/bin/desktop/force_run_policy.rs) +admits native carrier/particle material while refusing configured/session/ +provider/free-surface/AI workflows. Presence includes empty/disabled/zero scenes. +[FrozenSpec](../../../../src/bin/desktop/run_sets.rs#L248) preserves legacy +native intent v1 without fields and adds v2 with the validated scene and a +force-inclusive digest. This is frozen rerun intent, not accepted-state restart. +No native interaction or worker runtime was exercised by this refresh. + +## Proposed text and placement for root integration + +These locations refer to the merged original research text before publication +edits; headings remain stable if root changes line numbers. + +| Artifact/location | Required correction | +| --- | --- | +| CPU inventory introduction, currently lines 3–9 | Keep `fd7c0f0...` as original inspection provenance. Add a separate dated publication refresh at full `d67aec1...` and identify this report/current source. Do not globally replace the old pin and imply the old review executed on new source. | +| CPU inventory source map, section 1 | Add `force_field.rs`, its frozen scene/evaluator and force-aware recording owner. Scope the old constant-vector `BodyForce` paragraph as the legacy API, then add authored-field timing/support/density distinctions. | +| CPU inventory step order, sections 3 and 5 | Add authored bound and face force stage; name its closed-box-only combinations. Distinguish both invalidated carrier stamps from persistent accepted particle controls on rejection. | +| CPU inventory particle section 7 | Acceleration may include authored fields. Include local evaluated-density conversion and retry time semantics, while retaining pre-integration versus accepted-position timing. Keep MaterialSession's unchanged narrower API explicit. | +| CPU inventory recording section 8 | Qualify four/five fields as legacy, add one input document for authored profiles and run digest custody. Explain force metadata is typed while other text remains opaque. Preserve the no-restart conclusion. | +| Chapter 37 introduction and section 37.8 | Preserve original baseline, link dated refresh. Add actual field-evaluator semantics so radial/vortex are not confused with vorticity confinement or future buoyancy. Frozen scene/order/time/density are part of the admitted force kernel and dt reduction policy. | +| Chapter 38 source baseline and existing acceptance section | Identify the two source epochs. Add immutable force document/version/digest to device input lifetimes and readback admission; changing authoring draft must not mutate an in-flight device scene or relabel accepted records. These inputs do not alter CUDA/Vulkan independence. | +| Chapter 39 introduction and oracle tuple | Retain original pin and append current refresh; add exact force-scene bytes/version/digest, activation stage, density convention and recording profile to future oracle identity. | +| Chapter 39 validation matrix, Forces/Persistence rows | Add rotated radial/vortex/support edges, activation endpoints/straddles, empty-vs-absent, summed-object order, force-aware stamps and run/step scene mismatch. Core integration tests are prepared/ignored and were not executed here. | +| Chapter 39 G3 and following paragraph, currently lines 319–324 | Replace blanket exclusion of force recording with the exact legacy/authored/viscosity/GPU distinction below. | +| Book source plan, source baseline paragraph | Record PR145/146 as merged at the publication refresh, retaining that they were active/outside the original write set. Separate original research authorization/checks from later authorized publication; do not rewrite history as if navigation/build/merge had been authorized originally. | + +Suggested replacement for Chapter 39's recording-profile paragraph: + +> At the original research baseline, canonical carrier recording admitted the +> stamped `step` and `step_with_emitter` lineage. At the publication refresh +> `d67aec1a4575e0a61aa8cd931c8eeec33c8dc6a3`, it also consumes the force-aware +> success stamp from `step_with_force_fields` and +> `step_with_emitter_and_force_fields`, retaining the actual immutable scene in +> `external_force_scene_v1` and checking its run binding. Legacy `BodyForce` +> `step_with_forces`, viscosity, configured/provider/free-surface and GPU-resident +> producers are not thereby admitted. A new GPU adapter must preserve field +> stages, actual source/force custody, profile inventory and run/step association; +> passing existing recording tests does not admit a new producer. + +Suggested publication refresh text for the inventory and chapter/source-plan +introductions (abbreviate and link to the single detailed owner): + +> The original inventory is pinned to `fd7c0f0aee63bc8710df99f9bb01154e7619b974`. +> A source-only publication refresh inspected main +> `d67aec1a4575e0a61aa8cd931c8eeec33c8dc6a3` after PR145/146 merged. Pressure, +> layout and acceptance remain unchanged in that delta; immutable authored +> force scenes, their MAC/particle adapters and force-aware recording custody +> are additional current capabilities. This refresh adds no execution result. + +## GPU recommendation consequences and available source oracles + +Keep G0 plus a small isolated G1 dense operator lab as the first future slice. +An optional first no-authored-field profile must explicitly say that it excludes +the new force capability; a later full-step parity claim must either implement +that admitted profile or visibly refuse it. Do not silently drop nonzero fields, +map radial/vortex to a constant vector, normalize away authored bits or treat +empty metadata as absent. The GPU does not create new physics by evaluating +these existing formulas in parallel. + +For an authored MAC extension, retain one immutable, versioned/digested device +scene per admitted input epoch, evaluate on each own-axis face lattice at interval +start using constant rho, preserve the conservative preflight acceleration bound, +and preserve the f32 force work calculation. Define precision/order policy for +rotation, normalization, support smoothstep and object superposition. The +snapshot input and staging copies add memory beyond the old field-payload plan. +For particle work, use the separate local-density/pre-integration contract and +retain its retry/custody distinctions; it remains a later separate model. + +The relevant source oracles now include +[force-field contracts](../../../../tests/force_field_contract.rs), +[solver adapter contracts](../../../../tests/force_field_solver_contract.rs), +[force recording integration contracts](../../../../tests/force_recording_integration_contract.rs) +and in-source codec/schema adversarial tests. The two integration suites contain +explicit ignored execution admission; their presence or compilation cannot be +reported as numerical/SQLite acceptance. This refresh did not run, reinterpret +or authenticate their upstream receipts. Root's separately owned publication +checks must report their own exact scope and outcomes. diff --git a/docs/plans/gpu-grid-publication/reports/independent-publication-review.md b/docs/plans/gpu-grid-publication/reports/independent-publication-review.md new file mode 100644 index 00000000..d0850a65 --- /dev/null +++ b/docs/plans/gpu-grid-publication/reports/independent-publication-review.md @@ -0,0 +1,217 @@ +# Independent GPU publication source review + +Review date: 2026-10-21. Publication base: +`d67aec1a4575e0a61aa8cd931c8eeec33c8dc6a3`. Imported research snapshot: +`295e9be5736a92425b4281047576d959bbc1d5b2`. +Reviewer ownership: this report only. Chapter/shared-file repairs belong to the +publication owner. Book authorship remains **Puma**. + +**Final source disposition: no unresolved deciding technical, citation, license +or source-scope blocker in the frozen publication content reviewed below.** +The prose conversion, public audit and dated CPU refresh were reread. Committed +head identity, actual CI, PDF/browser results and deployment remain separate +subsequent evidence; this report does not predeclare them. +This is source/documentation review, not publication deployment or numerical, +device or performance qualification. + +## Applied publication rule + +The reviewer read the existing +[FEM publication source audit](../../../research-book/implementation/lagrangian-fem-sources-and-scope.md) +and [publication review](../../lagrangian-fem-book/reports/publication-review.md). +The applied rule preserves original mathematical exposition and scientific +provenance while excluding copyleft or uncertain implementation incorporation. +Scientific citation is distinct from reproducing copyrighted expression or +incorporating software. No third-party paper or proprietary API documentation +is thereby relicensed under Apache-2.0. No source, figure, substantial prose, +proof body, binary or dependency from the reviewed projects is imported. + +At the publication base, `LICENSE` blob +`261eeb9e9f8b2b4b0d119366dda99c6fd7d35c64` contains Apache License 2.0 and +`Cargo.toml` blob `7d6846d80bb3d98eace08eb6d6c042642592dbdc` declares +`Apache-2.0`. These match the earlier FEM audit; this publication changes neither. +The only implementation studies retained in the GPU chapters are the inspected +AmgX and modern Flow/PhysX pins, whose relevant license/source authorities are +BSD-3-Clause. This does not approve an entire dependency graph for future reuse. + +## Citation and license checks + +The original independent GPU source review remains historical evidence at its +declared CPU/research snapshot. For this publication review, the reviewer again +opened the following decisive primary authorities on 2026-10-21. + +| Authority | Inspected evidence and publication limit | +| --- | --- | +| [G7 Weber author manuscript](https://www.academia.edu/14777152/A_Cut_Cell_Geometric_Multigrid_Poisson_Solver_for_Fluid_Simulation) | Its author-copy header identifies the authors and volume 34(2), 2015; its abstract/body describes staggered finite-volume cut-cell hierarchy construction. Chapter 37 uses the narrow geometry-specific motivation. No rate, timing, detailed transfer formula or implementation guarantee is imported from it. Publisher/Fraunhofer failed access is retained as failed access, not full-text verification. | +| [G9 Aanjaneya author project abstract](https://orionquest.github.io/papers/EGMVL/paper.html) | Confirms authors/venue/year and coupled viscous stencils requiring box smoothing. Chapter 37's limited box-smoother statement is supported. The original strain-form equation and discussion are general exposition; no detailed paper algorithm, box-system size, Cholesky recipe, reported speedup or capacity is attributed to the unread article body. The bibliography explicitly limits evidence to the abstract. | +| [Pinned AmgX BSD license](https://raw.githubusercontent.com/NVIDIA/AMGX/91a8413ef267b1c32aff4014c02820e1c5897ac2/LICENSES/BSD-3-Clause.txt) | Reopened three-clause text, including notice/disclaimer and nonendorsement requirements. The separately inspected PCG/polynomial source SPDX headers and exact `91a8413...` pin bound the architectural comparison. No GPL implementation or dependency import is made. Bundled third-party licensing remains outside the narrow project-file conclusion. | +| [Pinned Flow/PhysX BSD license](https://raw.githubusercontent.com/NVIDIA-Omniverse/PhysX/b97365c6ddb32fda8469bc788ec18024862b0368/LICENSE.md) | Reopened BSD-3-Clause text at `b97365c...`. The inspected pressure sources carry matching SPDX headers. This is the modern PhysX Flow tree, not a license conclusion about older GameWorks Flow or every bundled component. The chapter correctly labels its centered collocated float4 divergence as different from Rheon's MAC pair. | + +The earlier independent review directly reopened the other deciding literature +and official API authorities: Harris/Crane GPU Gems, Bolz institutional PDF, +McAdams author PDF, Harris reduction presentation, Bridson/Müller-Fischer notes, +NVIDIA streams/graphs/platform/capability pages, and Khronos +features/atomics/shader/synchronization pages. Their provenance and use in this +publication remain subject to the same scoped facts: historical algorithms and +API rules support original design synthesis, without becoming Rheon benchmark, +installed-device, implementation, convergence or physical-validation evidence. + +The bibliography accurately distinguishes source retrieval from execution, +metadata from inspected bodies, exact repository hashes from moving vendor +documentation and source licensing from future combined-work licensing. In the +reviewer's initial session exact upstream GitHub commit timestamps were not +independently retrievable; the authors' recorded metadata check supplies those +timestamps. Pinned raw content/license identity was independently checked. + +## Deciding observations and final disposition + +No high- or medium-severity citation/license defect was found in the initial +publication drafts. The reviewer sent the owner these publication requirements +and boundaries for the final pass: + +1. The final public source audit should expressly separate factual paper/API + provenance from permissive implementation comparisons and state that no + copyrighted source expression or GPL software is imported. This follows the + applied FEM publication rule; it is not a blanket license judgment about + scientific citations or all potential future uses. **Resolved:** the new + [public audit](../../../research-book/implementation/gpu-sources-and-publication-scope.md), + lines 30–68, states these distinctions, exact BSD pins, no copied software or + substantial source expression and no GPL/copy-left dependency selection. +2. The retained CPU inventory freezes `fd7c0f0`; it must not be presented as a + complete current-main audit. The publication base has later force-field, + recording, particle and desktop changes. The pressure/operator/geometry + source owners are unchanged between those commits. The original frozen + oracle remains a sound declared comparison boundary. **Resolved:** the + inventory's dated publication refresh at lines 20–95 explicitly retains + sections 1–10 at the old source epoch, while describing current authored + force evaluation, adapter density/time, custody and typed recording scope. + Chapters 37–39 and the source plan link that distinction. +3. Root is removing the research snapshot's pseudocode/code-like diagrams. + Final review must verify that the pressure acceptance, restart direction, + final-budget classification, stored-velocity gate and publication ownership + survive as prose/equations, and that no fenced code block remains in the + three publication chapters. **Resolved:** all nine research code/diagram + blocks are replaced. Chapter 37 lines 138–163 retains the checked PCG + recurrence as prose and equations, including zero-problem acceptance, + positive finite products, fresh true residual, immediate final-budget + exhaustion and a per-iteration replacement decision. Pressure acceptance + still precedes f32 correction and actual stored-field divergence. Chapter + 38 retains independent API custody and generation/lease-based publication. + Chapter 39 retains allowed gauge alignment, acceptance and failure gates. + +## Final source and document checks + +The owner froze content and requested this final pass. The reviewer inspected +the changes relative to research commit `295e9be...` and publication base +`d67aec1...`: Chapters 37–39, the CPU inventory's dated refresh, the updated +reading/source plan, the new public publication audit, and the single additive +book README paragraph. The two GPU bibliographies remain unchanged from the +reviewed research snapshot. Retained historical plan/review files remain +historical and are not substituted as the current public source audit. No +production, dependency manifest, proof, admission, tolerance, workflow or book +builder change is present in this source scope. + +Independent Pandoc JSON parsing used `markdown+tex_math_single_backslash` and +compared each original Math node as a multiset, preserving exact expression +text and display/inline association. All **25 original expressions** remain; +there are **42 current expressions** after explanatory additions. The three +chapter ASTs contain **zero CodeBlock nodes**, compared with nine in the +research snapshot. Counts by chapter are: + +| Chapter | Original math | Current math | Missing original math | Current code blocks | +| --- | ---: | ---: | ---: | ---: | +| 37 | 10 | 24 | 0 | 0 | +| 38 | 0 | 0 | 0 | 0 | +| 39 | 15 | 18 | 0 | 0 | + +A separate local relative-link existence scan covered the three chapters, +three GPU implementation guides and two GPU bibliographies: **181 links, zero +missing destinations**. This check does not qualify fragment anchors, public +URLs, mathematical rendering, table pagination or browser/PDF results. +Staged and unstaged Git whitespace checks also passed. Root owns its separately +reported strict KaTeX, anchor, document-fixture and eventual pipeline checks. + +The reviewed content hashes below permit a later committed-head identity check +without a self-referential report hash. They identify source content, not an +artifact/deployment receipt. + +| Reviewed path relative to repository | SHA256 | +| --- | --- | +| `docs/research-book/chapters/37-gpu-grid-algorithms.md` | `21475f9a70436a9faa6a3497b26b87ab447928be61331219661afa6671e81120` | +| `docs/research-book/chapters/38-independent-cuda-and-vulkan.md` | `ad4139a764e35477b6b81450abccbec1b48a70eb59df37b3f14a9a41bbc3e275` | +| `docs/research-book/chapters/39-gpu-validation-and-roadmap.md` | `03b619ba8f8dab0e194d59794137599fdefe4a3883c7e423b4cb86cba4d06ec8` | +| `docs/research-book/implementation/gpu-cpu-oracle-inventory.md` | `72c9c9a2d39a63a75ec5e25b5ac549740cc32fc610078623819b1c641a8c104f` | +| `docs/research-book/implementation/gpu-research-plan.md` | `e95f2b1c4505317616de096747562215da4f4c6d2170c5aa43fb79c62ad8d384` | +| `docs/research-book/implementation/gpu-sources-and-publication-scope.md` | `c536c28145503698c5d3545ede1d4621fbde5c245707256de6ff3d251344aa8c` | +| `docs/research-book/sources/gpu-algorithm-bibliography.md` | `ce9ca2e928f30ef2e5e0c5cdf266cb0778d5ee6f0080e243625b08026c6bf509` | +| `docs/research-book/sources/gpu-backend-bibliography.md` | `8f9c10e6be93f19d41c3b2d5a63d6817c6860ca563bd74441efe95c50ea2f491` | +| `docs/research-book/README.md` | `d9581edf7218857166ea99842d9c3621b356e7049754601a91726127aab6ae6b` | + +## Evidence boundary and remaining gates + +No source-review action was denied or left with an uncertain mutation outcome. +The metadata access limit above remains a disclosed reviewer limitation, not a +claim of independently checked timestamps. There is no remaining source-content +blocker. A later exact-head identity review and the parent-owned normal +CI/PDF/browser gates are still required before merge/publication acceptance. +No solver/scientific fixture, GPU workload, hardware enumeration, workflow, +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/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/README.md b/docs/research-book/README.md index dde4d1c9..b374cdf6 100644 --- a/docs/research-book/README.md +++ b/docs/research-book/README.md @@ -4,6 +4,8 @@ A research book credited Puma: mathematics, algorithm specifications, worked exa The dedicated [Lagrangian FEM research expansion](implementation/lagrangian-fem-sources-and-scope.md) adds chapters [30](chapters/30-material-motion-and-conservation.md)–[36](chapters/36-reference-fem-solver.md): material motion and conservation, stress/weak forms, element operators, time/nonlinear solution, compressible and mixed fluids, contact/ALE and two fully specified minimal reference algorithms. The [publication source and scope audit](implementation/lagrangian-fem-sources-and-scope.md) records the solid/gas choices, license exclusions and remaining qualification. Chapters30–36 use prose, formulas, conditional proofs and worked examples without pseudocode. This is a reviewed draft exposition, not a general FEM implementation or new scientific qualification. The audit separates manuscript review from future numerical evidence. Puma's author credit and earlier research remain intact. +The GPU research expansion adds Chapters [37](chapters/37-gpu-grid-algorithms.md)–[39](chapters/39-gpu-validation-and-roadmap.md): compatible grid algorithms, independent CUDA and Vulkan compute approaches, and CPU-reference validation, device residency and staged implementation requirements. The [source/license and publication-scope audit](implementation/gpu-sources-and-publication-scope.md) preserves inspected-source/access limits and distinguishes original mathematical exposition from software reuse. This is prose-and-equation research without pseudocode, GPU implementation, hardware qualification or measured speedup; Puma's credit and existing numerical boundaries remain intact. + The [fluid portals and open boundaries chapter](chapters/40-fluid-portals-and-open-boundaries.md) adds Chapter 40, reserving 37–39 for separate GPU research. It develops periodic topology, hydrostatic reservoirs, weak/MAC conditions, conditional conservation and energy proofs, backflow and particle requirements. The [source/scope guide](implementation/fluid-portals-sources-and-scope.md) and [implementation proposal](../plans/fluid-portals-open-boundaries/plan.md) separate reviewed mathematics, publication evidence and future numerical acceptance. No portal implementation or new scientific qualification is claimed. The repository contains bounded Rust box/liquid components and separate research prototypes. General moving mesh/interface composition, coupled physical boundaries, particles and GPU sections remain incomplete or scoped research specifications. Exact proofs do not establish floating-point accuracy, geometry assembly, solver convergence, physical calibration or performance. 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..e912c353 --- /dev/null +++ b/docs/research-book/chapters/37-gpu-grid-algorithms.md @@ -0,0 +1,332 @@ +# 37 GPU algorithms for a compatible grid fluid + +*Research chapter for the Puma-credited Rheon book. Original CPU baseline: `fd7c0f0`; the [dated CPU refresh](../implementation/gpu-cpu-oracle-inventory.md) inspects publication-base `d67aec1` and authored-force/recording additions. The [publication source/license audit](../implementation/gpu-sources-and-publication-scope.md) preserves both source epochs and access limits. 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 explanation names dependencies without prescribing commands in a particular API. + +The interval starts from immutable accepted fields and a geometry snapshot. Time-step diagnostics precede velocity backtracing; force evaluation and any admitted viscous action consume that transported candidate. The resulting boundary flux defines the pressure right-hand side. Only after pressure qualification may face correction and stored-velocity divergence/Courant checks complete. Tracer or interface transport then consumes its declared carrier, and complete finite-value, bound, resource and cancellation gates precede publication of one accepted generation. Every dependency requires producer completion and visibility before its consumer reads that field version. + +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 gives each row one output owner: + +For each active pressure row, one invocation gathers the neighboring values and the authoritative face coefficients in fixed axis/side order. A fluid neighbor contributes \(t_{ij}(x_i-x_j)\); a known-pressure boundary contributes \(t_e x_i\) to the operator action, with its prescribed value already represented on the right-hand side. A prescribed-flux wall introduces no pressure coupling. The invocation writes that row's complete result exactly once. Input and output arrays must not alias, and all output writes must become visible before any consuming reduction or stencil begins. + +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..4f532828 --- /dev/null +++ b/docs/research-book/chapters/38-independent-cuda-and-vulkan.md @@ -0,0 +1,467 @@ +# 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. + +Original repository observations refer to `fd7c0f0`; a dated publication refresh +at `d67aec1` adds authored forces and recording custody while preserving +pressure/operator/geometry source. The +[CPU oracle inventory](../implementation/gpu-cpu-oracle-inventory.md) owns the +complete inventory; this chapter only extracts constraints that affect backend +design. The [publication source/license audit](../implementation/gpu-sources-and-publication-scope.md) +retains both source epochs and the licensed-comparison boundary. [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. + +The numerical contract and shared test inputs feed two separate implementations. The CUDA side owns CUDA capability discovery, allocations, modules, kernels, streams, events and optional graphs. The Vulkan side independently owns Vulkan capability discovery, buffers, device memory, SPIR-V, compute pipelines, queues and barriers. Both produce evidence for the same accepted-generation semantics; neither sits above or below the other in an execution adapter stack. + +This separation is a dependency design. 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: + +Capability probing establishes whether the requested numerical and resource policy is supported before creating a session. A session owns the admitted problem, accepted generation and memory budget. Submitting an attempt produces a ticket whose bounded status is pending, failed with a reason, or ready with a complete report. Readiness permits a separate generation-checked commit decision; it does not itself publish. Cancellation requests and bounded exports remain owned operations with their own completion obligations. + +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. + +For authored-force profiles, device input includes the exact frozen scene, +document version and digest admitted by the numerical owner. Support, activation +time and superposition order cannot change during an in-flight attempt. +Recording consumes actual accepted scene custody and verifies run/step binding; +a later authoring draft cannot relabel accepted device results. These inputs +and staging copies add to the budget without coupling CUDA to Vulkan. + +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 + +Initial upload completes before the explicitly selected compute stream reads any field or geometry. The stream transports accepted velocity into candidate storage, applies admitted forces and boundary treatment, assembles the right-hand side, and initializes pressure/Krylov state. Bounded chunks then apply the operator, reduce curvature, finalize checked scalar coefficients, update vectors, apply the preconditioner and reduce the next scalar products. Scheduled true-residual evaluation qualifies pressure. Every later dispatch in a completed chunk honors the current convergence or breakdown state without modifying a terminated solve. + +Pressure success precedes f32 velocity correction, actual stored-divergence and Courant checks, and tracer transport through its declared carrier. Remaining acceptance diagnostics form a bounded device report. A completion event makes that report available to a readback stream, which waits before copying into a pinned host slot and records its own copy-completion event. Only after observing that completion may the host inspect API status, attempt identity and cancellation and commit or reject the candidate. + +The described stages have 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 table 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 described below through its resource scopes. The backend must enable `synchronization2` and use a core +or extension entry point appropriate to the selected runtime profile. + +After operator application writes the product buffer, a same-queue synchronization2 buffer memory dependency covers the actual written range. Its source scope is compute-shader write and its destination scope is compute-shader read; queue-family indices are ignored for this same-family operation. Dot-product partials may then read that product. Each reduction level similarly makes its partial writes visible to the next level. Scalar-finalization writes require their own compute-write-to-compute-read dependency before the vector update reads alpha, beta or termination state. These dependencies must be expressed through the appropriate enabled core or extension entry point, with fully initialized barrier structures and correct offsets/ranges. + +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. + +An attempt reserves a candidate, scratch and a bounded report slot while keeping the accepted generation immutable. Before each bounded submission and at declared completion checkpoints, the owner observes cancellation and stops scheduling new work when requested. Already submitted work is retired before its resources can be reused. Correction, transport and all acceptance checks finish before the owner inspects the completed bounded report; any API error, nonfinite value or failed gate rejects publication. + +At the final publication decision, the owner synchronizes host-side custody and rechecks the observed cancellation state, session identity, buffer epoch and expected accepted generation. It then publishes the complete candidate fields, time, generation and report as one logical transaction. Old accepted storage remains leased until every renderer/export reader has completed; it cannot immediately become scratch merely because the numerical commit succeeded. + +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..951d0d31 --- /dev/null +++ b/docs/research-book/chapters/39-gpu-validation-and-roadmap.md @@ -0,0 +1,353 @@ +# 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. Its dated +publication refresh inspects `d67aec1` after authored-force/recording additions; +the [source/license audit](../implementation/gpu-sources-and-publication-scope.md) +keeps source epochs and publication qualification separate. + +## 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. +Include exact immutable force-scene bytes/version/digest, interval-start +activation, density convention and recording profile. 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 description specifies numerical and publication order without declaring a new runtime interface: + +One interval reads immutable accepted generation \(n\), its interval inputs and the admitted capability policy. A backend-owned candidate passes through velocity advection, forces and any separately admitted viscosity. Its right-hand side feeds pressure solution and full true-residual/predicted-divergence qualification. Velocity correction then precedes actual stored-divergence and interval gates. Tracer transport/source and all complete-step checks must succeed before generation \(n+1\) is committed. Bounded immutable snapshots are downstream recording/render inputs. An invalid, cancelled or unsuccessful candidate leaves generation \(n\) accepted. + +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/directional, rotated radial/vortex and support edges; force density with the admitted density convention | Units, own-axis sampling, activation endpoints/straddles, ordered superposition, source saturation, forbidden boundary DOFs; invalid 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, actual force custody and controlled writer failures | Field stages/units/identity, empty-versus-absent metadata, run/step scene mismatch, immutable render view, durable prefix/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 | + +At the original baseline, canonical carrier recording admitted stamped `step` +and `step_with_emitter` lineage. At the `d67aec1` refresh it also consumes actual +force-aware custody from `step_with_force_fields` and +`step_with_emitter_and_force_fields`, retaining the immutable scene through +`external_force_scene_v1` and its run binding. Legacy `BodyForce` +`step_with_forces`, viscosity, configured/provider/free-surface and GPU-resident +producers are not thereby admitted. A GPU adapter needs its own qualification +of stages, actual source/force custody, inventory and run/step association; +passing existing recording tests does not admit a new 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..431b2a17 --- /dev/null +++ b/docs/research-book/implementation/gpu-cpu-oracle-inventory.md @@ -0,0 +1,505 @@ +# 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. + +## Publication refresh: authored forces and recording + +Sections 1–10 below retain the original `fd7c0f0` source epoch. A separate +source-only refresh inspected publication-base main +`d67aec1a4575e0a61aa8cd931c8eeec33c8dc6a3` after PR145/146 merged. +The [publication source/license audit](gpu-sources-and-publication-scope.md) +keeps this refresh, historical execution and future GPU qualification separate. +Geometry, sampling/advection, pressure/operator, SGS, viscosity, liquid-step, +provider and session source are unchanged in that delta. No new numerical or +runtime pass is inferred from this comparison. + +[`force_field.rs`](../../../src/force_field.rs) adds version-one `ForceScene` +authoring and validated immutable `ForceSnapshot` evaluation for up to 256 +uniquely identified objects. Directional, signed constant-magnitude radial and +local-+Y vortex fields have world translation/quaternion orientation and global, +local open-sphere or local half-open-box support, with optional smoothstep edge +weight. These prescribed normalized fields are not inverse-square gravity or +solved vorticity confinement. Object order fixes superposition order. Activation +uses start-inclusive/end-exclusive bounds at the accepted interval start, with +no interval split at an activation endpoint. Disabled objects are validated; +empty/disabled/zero metadata is not silently converted to absent metadata. + +The new closed-box `Simulation::step_with_force_fields` and +`step_with_emitter_and_force_fields` do not compose tracer barriers, prescribed +flux, surface modes, viscosity, providers or legacy `BodyForce` slices. The +[MAC adapter](../../../src/forces.rs) evaluates each component at its own world +face position after transport and legacy smoke/nozzle forcing, before pressure. +Acceleration is m/s²; force density is N/m³ divided by constant carrier density. +Normal outer faces remain unchanged. Preflight uses the active-object magnitude +triangle bound including its rotation/rounding allowance, rather than only a +maximum sampled on occupied faces. Work uses actual stored f32 velocity changes +and discrete face mass. An empty scene skips the numerical field stage but +retains force-aware publication custody. + +The [particle adapter](../../../src/particle_fluid.rs) instead evaluates after +candidate emission and density/EOS, at pre-integration positions, before existing +pair-force accumulation. Force density divides by evaluated local SPH density, +not MAC density, reference density or parcel mass. All stability retries use the +accepted interval start. Diagnostics precede integration/contact while accepted +positions/velocities belong to the later publication time. The unchanged +`MaterialSession` wrapper does not expose every new direct API. + +Both carrier interval-pressure stamps, legacy and force-aware, invalidate on a +later prepare attempt even if it fails; accepted velocity/tracer/time/generation +remain intact. Accepted particle force controls own independent diagnostic +storage and survive rejected/cancelled attempts. A successful legacy particle +step clears authored custody; a successful authored step replaces it. These +owners must not be conflated in a GPU publication policy. + +[`recording/force_fields.rs`](../../../src/recording/force_fields.rs) adds the +typed little-endian `external_force_scene_v1` input document inside the unchanged +SQLite envelope. Capped at 256 KiB, it preserves order, IDs, names, disabled +objects and f64 bits including signed zero. It is immutable input metadata, not +sampled endpoint acceleration. Decoding bounds data and freezes the scene. +A run-header SHA256 binds the scene; append and selected read require matching +run/step metadata. Absence denotes the legacy profile, not proof that arbitrary +producers applied no force. + +The [carrier adapter](../../../src/recording/carrier.rs) now consumes either +actual legacy or force-aware success custody. Authored publication adds one input +field to the legacy four-constructor/five-accepted-field profile; its wrapper +cannot relabel an unstamped generation. The [material adapter](../../../src/recording/material.rs) +checks requested-dt bits, emitter and report against actual accepted force +custody. Most other controls remain opaque text, but this scene has a typed +validated codec. It and frozen rerun intent still omit complete restart state. +Durable-prefix, unrecorded-tail, uncertain-append poisoning and keep-all limits +remain unchanged. Legacy `step_with_forces`, viscosity and GPU producers gain +no recording admission simply because authored adapters exist. + +Current source oracles include [field](../../../tests/force_field_contract.rs), +[solver-adapter](../../../tests/force_field_solver_contract.rs) and +[recording](../../../tests/force_recording_integration_contract.rs) contracts. +Ignored integration cases are not fresh execution merely because they exist or +compile. Future parity identifies exact scene bytes/version/digest, evaluation +time/density, order/arithmetic and recording profile, or visibly refuses that +capability. No source review changes earlier physical-validation failures. + +## 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](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 + +[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](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](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 +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..0c337537 --- /dev/null +++ b/docs/research-book/implementation/gpu-research-plan.md @@ -0,0 +1,123 @@ +# 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 [publication source/license audit](gpu-sources-and-publication-scope.md) +owns the public source/access and publication-scope explanation. 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. Product PR145/146 were then active +and outside the original write set. Publication preparation now starts from main +`d67aec1a4575e0a61aa8cd931c8eeec33c8dc6a3`, after those product changes, FEM +PR147 and portal PR148 merged. The dated CPU inventory refresh records actual +authored-field and force-recording additions while preserving the original pin. +The publication changes no product source, proof, numerical gate, shared book +builder, workflow or existing bibliography; one additive book README paragraph +provides GPU discovery. Portal 40 and FEM 30–36 remain unchanged. + +## 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 were authorized +and executed in the original research phase. Later explicit publication +authorization admits the normal book/CI/PDF/browser gates, reported at their +actual source head separately from scientific GPU qualification. No GPU/hardware +workload, GPU backend implementation or benchmark is admitted. No generated +numerical artifact or copied paper is committed. Existing author credit remains Puma. + +Parent integration owns the final merge/publication decision after exact-head +CI and source/PDF/browser evidence. Nine earlier code/diagram blocks have been +replaced with prose and equations; original formulas and algorithm conditions +remain. Markdown tables 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/implementation/gpu-sources-and-publication-scope.md b/docs/research-book/implementation/gpu-sources-and-publication-scope.md new file mode 100644 index 00000000..c7d37cfc --- /dev/null +++ b/docs/research-book/implementation/gpu-sources-and-publication-scope.md @@ -0,0 +1,128 @@ +# GPU research: publication sources, licenses and evidence scope + +Book author: **Puma**. Publication preparation: 21 October 2026. This guide +accompanies Chapters [37](../chapters/37-gpu-grid-algorithms.md), +[38](../chapters/38-independent-cuda-and-vulkan.md) and +[39](../chapters/39-gpu-validation-and-roadmap.md). Their prose, equations and +derivations are original exposition. The publication replaces the earlier +research draft's nine code/diagram blocks with explanatory prose and explicit +recurrences; it contains no pseudocode, copied implementation code or new GPU +backend. CUDA and Vulkan remain independently usable proposed compute paths. + +## Preserved research and refreshed CPU scope + +Original research commit `295e9be5736a92425b4281047576d959bbc1d5b2` remains +preserved in the publication branch's ancestry. Its CPU oracle inventory pins +`fd7c0f0aee63bc8710df99f9bb01154e7619b974`, not an unspecified latest solver. +Publication preparation is based on main +`d67aec1a4575e0a61aa8cd931c8eeec33c8dc6a3`, after FEM PR147, force/backend +PR145, scene UI PR146 and portal PR148. No production physics, force eligibility, +recording schema, proof, tolerance or admission manifest changes are made here. + +The [CPU inventory](gpu-cpu-oracle-inventory.md) retains the original numerical +baseline and adds a dated force/recording refresh. The pressure/operator/geometry +files are unchanged across these bases; authored immutable force scenes and +their MAC/particle/recording paths are additional admitted CPU capabilities. +Those additions do not establish a device-resident interface or a general +free-surface/coupled-fluid solver. The [reading/source map](gpu-research-plan.md) +connects stages, model limits and proposed implementation gates. + +## License-compatible implementation references and scientific provenance + +Rheon's [LICENSE](https://github.com/MrScripty/Rheon/blob/d67aec1a4575e0a61aa8cd931c8eeec33c8dc6a3/LICENSE), +blob `261eeb9e9f8b2b4b0d119366dda99c6fd7d35c64`, contains Apache License 2.0. +[Cargo.toml](https://github.com/MrScripty/Rheon/blob/d67aec1a4575e0a61aa8cd931c8eeec33c8dc6a3/Cargo.toml), +blob `7d6846d80bb3d98eace08eb6d6c042642592dbdc`, declares `Apache-2.0`. +Neither file is changed by this publication. + +The user's publication rule excludes incompatible or unestablished software +reuse. The retained implementation comparisons are the exact AmgX and modern +Flow/PhysX source snapshots in the [algorithm bibliography](../sources/gpu-algorithm-bibliography.md). +Their inspected license authorities and selected source headers are BSD-3-Clause: + +- AmgX commit `91a8413ef267b1c32aff4014c02820e1c5897ac2`, version-file value + 2.5.0: [license authority](https://github.com/NVIDIA/AMGX/blob/91a8413ef267b1c32aff4014c02820e1c5897ac2/LICENSES/BSD-3-Clause.txt). + Selected PCG, polynomial smoother and configuration source supports factual + algorithm-composition comparisons. It is a CUDA library, not a Vulkan backend + or a complete Rheon fluid solver. +- Flow within PhysX commit `b97365c6ddb32fda8469bc788ec18024862b0368`: + [root license](https://github.com/NVIDIA-Omniverse/PhysX/blob/b97365c6ddb32fda8469bc788ec18024862b0368/LICENSE.md) + and selected pressure-source SPDX headers. These support sparse pressure-pass + orchestration and collocated-shader observations. This is not the older Flow + repository's unestablished license and not a MAC-compatible drop-in. + +No external code, substantial prose, figures, binaries, libraries, SDKs or +datasets are incorporated. No GPL/copy-left implementation dependency is selected +or imported. BSD notice, disclaimer and nonendorsement obligations would matter +to actual redistribution; this source comparison does not admit an entire +dependency graph or certify every bundled component. Future code adoption needs +the exact component/version, notices, transitives and distribution audit. + +Original papers, course notes, NVIDIA technical chapters and official CUDA and +Vulkan documents supply factual and scientific provenance. They are not being +redistributed or relicensed under Apache-2.0. Reading a mathematical identity or +API requirement is distinct from copying copyrighted expression or incorporating +licensed software. No vendor code sample, paper illustration, proof body or +substantial quotation is reproduced. Local recurrences, error metrics, resource +models, schedules and architecture recommendations are the book's original +synthesis under their stated assumptions. + +## Exact supporting claims and access limits + +The [algorithm bibliography](../sources/gpu-algorithm-bibliography.md) owns exact +paper identities, versions, dates, selected source paths and retrieval limits. +The [backend bibliography](../sources/gpu-backend-bibliography.md) owns official +CUDA/Khronos sections and moving-document revision limits. Access was recorded +on 2026-10-21; it is not an installed SDK or tested driver identity. + +| Authority | Admitted use and limit | +| --- | --- | +| Harris GPU Gems 38 (2004), G1 | Historical rectangular 2D grid-pass decomposition; no free-surface or current Rheon speed claim | +| Crane/Llamas/Tariq GPU Gems 3 chapter 30 (2007), G2 | Historical 3D pass/advection exposition and its collocated/MAC distinction; relevant sections, not a claim of whole-method parity | +| Bolz et al. (2003), G3 | Primary GPU sparse-CG/regular-multigrid formulation and reduction dependencies; historical hardware and stopping conditions | +| McAdams/Sifakis/Teran (2010), G4 | Boundary-aware multigrid-PCG ideas; the inspected implementation/timings are shared-memory CPU evidence | +| Harris reduction presentation, G5 | Hierarchical reduction construction; no invented publication date or modern API synchronization authority | +| Bridson/Müller-Fischer notes (2007), G6 | Pressure, MIC and boundary derivations in inspected sections; not a GPU triangular-solve implementation | +| Weber et al. (2015), G7 | Inspected author-manuscript motivation for geometry-specific cut-cell hierarchy; publisher full text unavailable, no imported rate/timing or general-geometry guarantee | +| Aanjaneya et al. (2019), G9 | Author-abstract-level coupled-viscous/box-smoother point only; full paper unavailable within retrieval limits, no detailed smoother reconstruction or timing claim | +| Official NVIDIA guide/runtime sources | Selected execution, memory, graph, arithmetic and capability requirements; no entire SDK audit, hardware query or supported-model certification | +| Official Khronos specification/reference pages | Selected compute, resource, feature, subgroup and synchronization requirements; not a whole-specification read or a requirement that every Vulkan device supports f64 | + +Relevant sections of G2, G5, G6 and vendor documents were consulted, rather than +asserted cover-to-cover readings. G7 uses author-manuscript text, not a successful +publisher full-text retrieval. Exact AmgX/Flow commit timestamps were checked by +the bibliography author; the independent reviewer reopened pinned source and +licenses but could not reproduce that metadata request. These limits are retained +in the public ledgers rather than silently promoted to full verification. + +The consulted Vulkan documentation identifies build 1.4.365, generated +2026-10-02 from `8c9361ba8180c1f4164c0bf79de2f6e817770b0d`. Checked CUDA Guide +footers identify 2026-09-10; the separately consulted Runtime API page identifies +13.4. These moving references do not identify an installed toolkit or set a +minimum runtime profile. Optional FP64, atomic and subgroup behavior must be +queried and enabled for each later implementation/device configuration. + +## Publication evidence is separate from solver qualification + +The original independent source review repaired pressure/stored-field acceptance +ordering, per-iteration replacement state, final-budget failure classification, +gauge comparison and a chapter link. The publication continuation independently +reviews the prose-only recurrences, updated CPU/source scope and license/access +limits. Final commit, exact-head CI, actual PDF visual inspection, source/link +and browser evidence belong to the publication completion report; this guide +does not predeclare those results. + +The existing builder discovers all three chapters and this implementation guide. +Chapter bodies enter the normal reading/print/PDF edition. This audit is a +separate HTML implementation page and packaged Markdown source; its full body +is not automatically appended to print/PDF because that sequence uses the +existing explicit implementation-guide list. No fabricated interactive GPU lab +or benchmark is added. Existing FEM 30–36, portal 40 and thermal-owned paths are +outside this task's write scope. + +A successful book pipeline qualifies its actual source/artifact and publication +checks. It cannot establish an implemented GPU solver, numerical CPU/GPU parity, +physical convergence, hardware support, restart equivalence or measured speedup. +Those remain prospective, separately authorized gates in Chapter 39. The earlier +Library preview and uncertain upload state are retained unchanged; repository +publication is a distinct operation and does not retry that upload. 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.