wip(standalone): GPU step 7 -- viewer cold start 2.45 -> 1.49 s on CPU; GPU skirts missing

Resumed round: checkpoint rebuilt and re-verified (parity 3,436/3,436, walk
5,642/5,642, W1/W8 hashes 44/44, far topology 0 diffs vs CPU and WARP, far
survival PASS). Remote GPU holes were missing clipped LOD skirts: tiles with
visible skirt faces now fall back to the CPU mesh path, leaving only 16 far
tiles GPU-rendered until a GPU skirt emitter exists.

My interleaved runs (SkirtSeamVerify, process start -> whole view):
CPU 1.522/1.488/1.430 s; GPU 4.271/1.749/1.659 s. A ~2.8 s first Present
stall hit only my first GPU launch; Luna saw it on every fresh launch, so the
trigger is still unidentified. Headless CPU clipmap 0.789 s.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MmSoF64BbRFmk9C3p1VsBJ

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep7CaptureMatrix.ps1

@@ -0,0 +1,68 @@
+[CmdletBinding()]
+param([string]$ExecutableName='VoxelForgeGpuFieldStep7SkirtSeamVerify.exe')
+
+$ErrorActionPreference='Stop'
+$HostProcess=Get-Process -Id $PID
+$HostProcess.PriorityClass='BelowNormal'
+$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$ExecutablePath=Join-Path $StandaloneRoot ('bin\'+$ExecutableName)
+if([IO.Path]::GetFileName($ExecutableName)-ne$ExecutableName-or[IO.Path]::GetExtension($ExecutableName)-ne'.exe')
+{throw 'ExecutableName must be an executable filename.'}
+if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "Viewer not found: $ExecutablePath"}
+$Views=@(
+    @{Name='chamber';Camera='-11.75,-123.25,-19.85,0,0'},
+    @{Name='tunnel';Camera='-11.75,-123.25,-19.85,90,0'},
+    @{Name='passage';Camera='-39.19,6.27,-28.51,63,0'},
+    @{Name='far';Camera='-11.61244,-118.2212,-19.8452,0,0'},
+    @{Name='passage_floor_closeup';Camera='-39.19,6.27,-25.51,63,-25'},
+    @{Name='tunnel_bend';Camera='-198.97,-99.57,-45.58,33.59,23.17'},
+    @{Name='remote_yaw0';Camera='500,0,0,0,0'},
+    @{Name='remote_yaw90';Camera='500,0,0,90,0'},
+    @{Name='remote_yaw180';Camera='500,0,0,180,0'},
+    @{Name='remote_yaw270';Camera='500,0,0,270,0'}
+)
+foreach($View in $Views)
+{
+    foreach($Mode in @('cpu','gpu'))
+    {
+        $Stem="step7_${Mode}_$($View.Name)"
+        $CaptureName="$Stem.png"
+        $Arguments=@('--clip-radius=3','--workers=8','--quit-after-clipmap','--capture-after-clipmap',
+            "--camera=$($View.Camera)","--$Mode-far",
+            "--capture=Research/OwnerShapeWorld/gpu_field_port/$CaptureName")
+        $StdoutPath=Join-Path $PSScriptRoot ($Stem+'.stdout.log')
+        $StderrPath=Join-Path $PSScriptRoot ($Stem+'.stderr.log')
+        $CapturePath=Join-Path $PSScriptRoot $CaptureName
+        Write-Output "CAPTURE_COMMAND mode=$Mode view=$($View.Name) `"$ExecutablePath`" $($Arguments -join ' ')"
+        for($Attempt=0;$Attempt-lt2;$Attempt++)
+        {
+            if($Attempt-gt0)
+            {Write-Host '0xC0000142 encountered; waiting 60 seconds before the single permitted retry.';Start-Sleep -Seconds 60}
+            Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
+            $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
+                -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
+            $PeakWorkingSet=0L;$PeakPrivateMemory=0L;$OverMemoryLimit=$false
+            while(-not$Child.HasExited)
+            {
+                $Child.Refresh();if($Child.HasExited){break}
+                $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+                $PeakWorkingSet=[Math]::Max($PeakWorkingSet,$Child.WorkingSet64)
+                $PeakPrivateMemory=[Math]::Max($PeakPrivateMemory,$Child.PrivateMemorySize64)
+                if($Child.PrivateMemorySize64-ge8GB)
+                {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+                Start-Sleep -Milliseconds 100
+            }
+            $Child.WaitForExit();$Child.Refresh()
+            if($Child.ExitCode-eq-1073741502-and$Attempt-eq0){continue}
+            if($Child.ExitCode-eq-1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
+            if($OverMemoryLimit){throw "Viewer stopped at the private-memory limit: $Stem"}
+            $LogText=Get-Content -LiteralPath $StdoutPath -Raw
+            if(-not$LogText.Contains('[InitialClipmapReady]')-or-not(Test-Path -LiteralPath $CapturePath -PathType Leaf))
+            {throw "Capture or clipmap readiness missing: $Stem, exit=$($Child.ExitCode)"}
+            Write-Output "CAPTURE_DONE mode=$Mode view=$($View.Name) pid=$($Child.Id) priority=BelowNormal affinity=0xFF00 exit=$($Child.ExitCode) peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) path=$CapturePath"
+            if($Child.ExitCode-ne0){throw "Viewer returned $($Child.ExitCode): $Stem"}
+            break
+        }
+    }
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/STEP7_RESUME_REPORT.md

@@ -0,0 +1,65 @@
+# Step 7 resume report
+
+Checkpoint: `41716f0`. No commits, pushes, or stashes were made. The only tracked edit is in `Standalone/Source/StandaloneViewer.cpp`; the CPU field and mesher sources remain unchanged.
+
+## Checkpoint verification
+
+Fresh builds succeeded as `VoxelForgeGpuFieldStep7ResumeVerifyTimeline.exe` and `VoxelForgeGpuFieldStep7ResumeVerifyRunner.exe`, using `/MP8`, BelowNormal priority, and affinity `0xFF00`.
+
+- `--parity-only`: 3,436/3,436 samples bit-identical, 0 mismatches; the deferred diagnostic measured 1,057.646 ms.
+- Headless CPU `--stream-benchmark --shape-world --force-generate --stream-hash --workers=8 --clip-radius=3`: 2,226/2,226 tiles in 0.789 s; walkability 5,642/5,642.
+- The matching one-worker run took 4.108 s and passed walkability 5,642/5,642. W1/W8 hashes match in all 44 fields, and W8 matches the saved speed-round reference in all 44 fields.
+
+## Remote holes and seam handling
+
+Both CPU and original GPU captures had all 2,226 tiles resident. The GPU image still showed openings, so this was not a capture taken before residency. The CPU mesh includes clipped LOD boundary skirts; the GPU path emitted marching-cubes surfaces and connector repairs but no skirts. On a representative LOD2 tile, the CPU core surface had 289 vertices and the complete CPU mesh had 481, with 192 additional skirt vertices. The GPU marching-cubes surface agreed with the CPU surface within the measured 0.05 cm bound.
+
+The viewer now routes tiles with visible skirt faces through the existing CPU mesh path. This closes the remote openings and preserves the CPU's clipped skirt geometry and face masks. It is a correctness fallback with a substantial performance cost: in the default view, 1,738 tiles needed visible skirts, leaving 336 GPU tiles total (320 near, 16 far); in the remote yaw-0 view, 1,810 tiles needed skirts, leaving 300 GPU tiles total (288 near, 12 far). A GPU skirt kernel is needed to restore GPU coverage across the far rings.
+
+The full far GPU agreement rerun completed on `step7_remote_far_audit.stdout.log`: 1,150 tested tiles, 0 topology-different cells against CPU and WARP, 145,494 protected cells and 88 thin-feature cells with 0 differences on either implementation, and bit-identical repeated GPU outputs. Route-air, seal, and thin-route samples were exact.
+
+The GPU-mesh far-survival audit passed on the full far-GPU path. It covered 7,200/7,200 strata boundary surfaces and 10,976/10,976 far connectors. Connector sample/ray misses and holes were zero; the audit also counted 50,019 extra connector ray hits. The audit log is `step7_gpu_far_survival_reaudit.stdout.log`.
+
+## First `Present(1)` stall
+
+Three fresh launches of the same GPU viewer binary had first-present calls of 3.248, 3.007, and 3.066 s. CPU-mode first presents were 1.568–1.941 ms. GPU shader-object creation took about 52–75 ms per launch, with precompiled bytecode. Later presents in the same runs were under 5.4 ms.
+
+The stall repeated after each new process launch, so a one-time driver-cache miss that disappears on subsequent launches does not explain it. A diagnostic `--no-vsync` run still blocked for 3.137 s on its first present, ruling out ordinary sync-interval waiting as the sole cause. GPU timestamped work took 34.6–44.6 ms in these runs while the first batch completion and present were delayed by roughly 3.1–3.5 s. The remaining delay is in first GPU-batch/device/queue startup; these measurements do not separate driver pipeline setup from command-queue synchronization.
+
+## Cold whole-view timings
+
+Command: `RunStep7ColdInterleaved.ps1 -ExecutableName VoxelForgeGpuFieldStep7SkirtSeamVerify.exe`. Sequence: CPU A, GPU A, GPU B, CPU B, CPU C, GPU C; each used `--clip-radius=3 --workers=8 --quit-after-clipmap` and the matching `--cpu-far` or `--gpu-far` mode.
+
+| Mode | Process start to full view | Median |
+|---|---:|---:|
+| CPU | 1.645, 1.473, 1.492 s | 1.492 s |
+| GPU hybrid with seam fallback | 4.611, 4.122, 4.258 s | 4.258 s |
+
+The current hybrid takes 2.85× the CPU median. The CPU headless full-clipmap generation clock is 0.789 s at eight workers. There is no whole-view headless GPU mode; the separate 1,150-tile far audit measured 111.450 ms of GPU field-plus-meshing, while its CPU `generateTile` reference sum was 2,080.238 ms. These are component timings, not a whole-view headless comparison.
+
+For the current 336 GPU tiles, the GPU timestamp total was 34.6–42.4 ms, averaging 0.103–0.126 ms per tile. Four 90 Hz slices average 8.6–10.6 ms of GPU work each, under the 11.11 ms frame budget. This excludes CPU fallback generation and input preparation; those remain the wall-clock bottleneck. GPU input preparation took 128–162 ms wall / 0.94–1.21 worker-seconds. The DXGI sampled local-memory peak was 672.6 MB in the interleaved run.
+
+The full-GPU survival audit's higher-readback run reached a 2.978 GB DXGI local-memory peak and 4.73 GB process-private peak. Their conservative sum is 7.71 GB, under the 8 GB limit. Viewer capture runs peaked at 2.01 GB process-private.
+
+## Full-residency captures
+
+`RunStep7CaptureMatrix.ps1` produced six fixed-view and four remote-view CPU/GPU pairs. Every run required `[InitialClipmapReady]` before accepting its 1280x800 capture. The numerical comparison is in [step7_capture_comparison.md](step7_capture_comparison.md).
+
+| View | MAE RGB |
+|---|---:|
+| Chamber | 0.6075 |
+| Tunnel | 0.6383 |
+| Passage | 1.6092 |
+| Far | 0.7910 |
+| Passage floor close-up | 0.9314 |
+| Tunnel bend | 0.8129 |
+| Remote yaw 0 | 0.0258 |
+| Remote yaw 90 | 0.0097 |
+| Remote yaw 180 | 0.0451 |
+| Remote yaw 270 | 0.0815 |
+
+The remote views no longer show the former large blue openings. The seam fallback makes those captures visually close to CPU, at the GPU-coverage cost above.
+
+## Current outstanding work
+
+The viewer-level seam fallback restores the expected remote look, but it leaves too few far tiles on the GPU for a whole-view speedup. A GPU skirt emitter must reproduce the CPU's face-filtered, field-clipped skirts before the GPU far path can retain both the captures and broad far-ring coverage. The per-process first GPU-batch stall also remains unexplained beyond ruling out ordinary VSync waiting and a cache miss that warms across process launches.

Standalone/Research/OwnerShapeWorld/gpu_field_port/step7_capture_comparison.md

@@ -0,0 +1,16 @@
+# Step 7 CPU/GPU capture comparison
+
+Each pair is a 1280x800 back-buffer capture from the fixed camera. `mae_rgb` and `rmse_rgb` are channel errors on the 0 to 255 scale; pixel percentiles use the maximum RGB channel error per pixel.
+
+| View | MAE RGB | RMSE RGB | p95 pixel max | p99 pixel max | Max channel | Changed pixels | >8 | >32 |
+|---|---:|---:|---:|---:|---:|---:|---:|---:|
+| chamber | 0.6075 | 1.8642 | 3 | 7 | 123 | 42.64% | 0.60% | 0.02% |
+| tunnel | 0.6383 | 3.3799 | 3 | 8 | 130 | 31.73% | 0.87% | 0.28% |
+| passage | 1.6092 | 3.6878 | 12 | 20 | 52 | 50.50% | 8.95% | 0.17% |
+| far | 0.7910 | 2.1578 | 5 | 9 | 128 | 44.83% | 1.25% | 0.02% |
+| passage_floor_closeup | 0.9314 | 2.4800 | 5 | 13 | 100 | 43.02% | 2.53% | 0.05% |
+| tunnel_bend | 0.8129 | 2.1599 | 5 | 11 | 55 | 40.96% | 1.80% | 0.07% |
+| remote_yaw0 | 0.0258 | 0.5435 | 0 | 1 | 90 | 2.71% | 0.01% | 0.01% |
+| remote_yaw90 | 0.0097 | 0.4230 | 0 | 0 | 94 | 0.96% | 0.01% | 0.00% |
+| remote_yaw180 | 0.0451 | 1.2945 | 0 | 1 | 102 | 2.87% | 0.04% | 0.03% |
+| remote_yaw270 | 0.0815 | 1.6867 | 1 | 1 | 107 | 5.01% | 0.08% | 0.06% |

Standalone/Source/StandaloneViewer.cpp

@@ -1486,6 +1486,7 @@ namespace
         bool bGpuFarEnabled=false;
         bool bGpuFarAuditEnabled=false;
         bool bGpuFarAuditReported=false;
+        bool bDisableVSync=false;
         std::unique_ptr<FFarSurvivalAudit> GpuFarAudit;
         bool bGpuFarFirstTileReported=false;
         std::filesystem::path GpuFarShaderPath;
@@ -2380,13 +2381,25 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld())return;
             std::vector<FVoxelTileKey> CpuTiles;CpuTiles.reserve(Keys.size());
             std::vector<FVoxelTileKey> Eligible;Eligible.reserve(Keys.size());
+            uint64 SkirtTransitionFallbackTiles=0;
             for(const FVoxelTileKey& Key:Keys)
             {
                 const VoxelForgeCore::TileKey CoreKey{{Key.Coord.X,Key.Coord.Y,Key.Coord.Z},Key.Level};
                 const bool NearGpuEligible=Key.Level<World->FullResClipLevels;
                 if(NearGpuEligible||ShapeWorld->isFarFieldTileEligible(CoreKey))
                 {
-                    Eligible.push_back(Key);
+                    const uint8 VisibleSkirtFaces=World->bGenerateSkirts
+                        ?ComputeVisibleSkirtFaceMask(Key,DesiredTileSet):0;
+                    if(VisibleSkirtFaces!=0)
+                    {
+                        // CPU tiles carry the established clipped LOD skirts. The GPU field
+                        // and marching-cubes outputs match the CPU surface, but currently do
+                        // not emit these seam curtains. Keep tiles with exposed LOD faces on
+                        // the reference mesh path until a GPU skirt kernel is available.
+                        CpuTiles.push_back(Key);
+                        ++SkirtTransitionFallbackTiles;
+                    }
+                    else Eligible.push_back(Key);
                 }
                 else CpuTiles.push_back(Key);
             }
@@ -2460,6 +2473,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     {return Key.Level<World->FullResClipLevels;}));
             std::printf("[GpuFieldEligibility] eligible_gpu_tiles=%zu near_gpu_tiles=%zu cpu_fallback_tiles=%zu\n",
                 GpuFarEligibleTiles.size(),NearGpuTiles,Keys.size()-GpuFarEligibleTiles.size());
+            std::printf("[GpuSkirtSeamFallback] tiles=%llu reason=visible_lod_skirt_faces status=active\n",
+                static_cast<unsigned long long>(SkirtTransitionFallbackTiles));
         }
 
         void PumpGpuFarGeneration()
@@ -3391,7 +3406,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             }
             CaptureRenderedFrame();
             const auto BeforePresent=std::chrono::steady_clock::now();
-            const UINT PresentSyncInterval=1u;
+            const UINT PresentSyncInterval=bDisableVSync?0u:1u;
             const auto PresentCallStart=std::chrono::steady_clock::now();
             SwapChain->Present(PresentSyncInterval,0);
             const auto PresentCallEnd=std::chrono::steady_clock::now();
@@ -5214,6 +5229,7 @@ int main(int Argc,char** Argv)
     {
         const std::string_view Arg=Argv[I];
         if(Arg.starts_with("--capture="))App.CapturePath=std::string(Arg.substr(10));
+        else if(Arg=="--no-vsync")App.bDisableVSync=true;
         else if(Arg=="--profile-attribution")App.bProfileAttribution=true;
         else if(Arg=="--profile-detail-band")bProfileDetailBand=true;
         else if(Arg=="--compare-evaluators")App.bCompareEvaluators=true;