wip(standalone): GPU step 15 -- cold-start drain: hybrid whole view 3.2 -> 1.25-1.68 s

Timeline showed the post-export window was frame-paced: the 4-batch / 2 ms
per-frame upload budget held CPU-fallback tiles up to 755 ms after they
finished, and 35 GPU batches drained one per frame. Before the first whole
view, GPU work now drains as one batch and CPU fallback starts with GPU input
preparation (--gpu-far-no-cold-drain = paced control).

Luna's race: CPU 1.393 s vs GPU 1.677 s. My interleaved runs (Step15Timeline,
RunGpuFarViewer): GPU 2.902 (first launch) / 1.252 / 1.209 s; CPU 2.196 /
2.250 / 2.061 s. The CPU viewer still uses the paced upload budget at cold
start (stream 674 -> 2250 ms in my runs vs 741 -> 1422 in Luna's), so the
race is not yet fair. Audits: hashes, walk 5,642/5,642, 0 diffs vs CPU/WARP.

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

Standalone/Research/OwnerShapeWorld/gpu_field_port/AnalyzeStep15Timeline.py

@@ -0,0 +1,152 @@
+#!/usr/bin/env python3
+"""Expand Step 15 timeline logs into per-frame, per-batch, and per-CPU-tile CSVs."""
+
+import argparse
+import csv
+import re
+from pathlib import Path
+
+
+def write_csv(path: Path, columns: list[str], rows: list[dict[str, object]]) -> None:
+    with path.open("w", newline="", encoding="utf-8") as stream:
+        writer = csv.DictWriter(stream, fieldnames=columns)
+        writer.writeheader()
+        writer.writerows(rows)
+
+
+def optional_float(match: re.Match[str] | None, group: str = 1) -> float | None:
+    return float(match.group(group)) if match else None
+
+
+def optional_int(match: re.Match[str] | None, group: str = 1) -> int | None:
+    return int(match.group(group)) if match else None
+
+
+def main() -> int:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("log", type=Path)
+    parser.add_argument("--prefix", type=Path)
+    args = parser.parse_args()
+    log_path = args.log.resolve()
+    prefix = args.prefix or log_path.with_suffix("")
+    if not prefix.is_absolute():
+        prefix = Path.cwd() / prefix
+
+    cpu: dict[str, dict[str, object]] = {}
+    gpu_batches: dict[int, dict[str, object]] = {}
+    frames: dict[int, dict[str, object]] = {}
+    cpu_event_re = re.compile(
+        r"^\[CpuTileTimeline\] event=(queued|started|finished|uploaded|resident) "
+        r"key=\(([^)]*)\) at_ms=([0-9.]+) frame=([0-9]+)"
+    )
+    gpu_batch_re = re.compile(
+        r"^\[GpuBatchTimeline\] event=(prepared|submitted|completed|uploaded|resident) "
+        r"batch=([0-9]+) at_ms=([0-9.]+) frame=([0-9]+)"
+    )
+    frame_re = re.compile(r"^\[ColdFrame\] stage=(begin|end) frame=([0-9]+) at_ms=([0-9.]+)")
+
+    for line in log_path.read_text(encoding="utf-8", errors="replace").splitlines():
+        match = cpu_event_re.search(line)
+        if match:
+            event, key, at_ms, frame = match.groups()
+            row = cpu.setdefault(key, {"key": key})
+            row[f"{event}_ms"] = float(at_ms)
+            row[f"{event}_frame"] = int(frame)
+            for field in ("queued", "started", "finished", "worker"):
+                field_match = re.search(rf"\b{field}_ms=([0-9.]+)", line)
+                if field_match:
+                    row[f"{field}_ms"] = float(field_match.group(1))
+            batch_match = re.search(r"\bbatch=(\(lod[^)]*\))", line)
+            if batch_match:
+                row["cpu_batch"] = batch_match.group(1)
+            continue
+
+        match = gpu_batch_re.search(line)
+        if match:
+            event, batch_id, at_ms, frame = match.groups()
+            row = gpu_batches.setdefault(int(batch_id), {"batch": int(batch_id)})
+            row[f"{event}_ms"] = float(at_ms)
+            row[f"{event}_frame"] = int(frame)
+            tile_count = re.search(r"\btile_count=([0-9]+)", line)
+            if tile_count:
+                row["tile_count"] = int(tile_count.group(1))
+            continue
+
+        match = frame_re.search(line)
+        if match:
+            stage, frame, at_ms = match.groups()
+            row = frames.setdefault(int(frame), {"frame": int(frame)})
+            row[f"{stage}_ms"] = float(at_ms)
+            for field in ("inter_draw_gap_ms", "cpu_pump_ms", "gpu_pump_ms", "pre_present_ms", "present_ms"):
+                field_match = re.search(rf"\b{field}=([0-9.]+)", line)
+                if field_match:
+                    row[field] = float(field_match.group(1))
+            present_sync = re.search(r"\bpresent_sync=([0-9]+)", line)
+            if present_sync:
+                row["present_sync"] = int(present_sync.group(1))
+            held = re.search(r"\bheld=([01])", line)
+            if held:
+                row["held"] = int(held.group(1))
+
+    cpu_rows = sorted(cpu.values(), key=lambda row: str(row["key"]))
+    for row in cpu_rows:
+        queued = row.get("queued_ms")
+        started = row.get("started_ms")
+        finished = row.get("finished_ms")
+        uploaded = row.get("uploaded_ms")
+        row["queue_wait_ms"] = started - queued if queued is not None and started is not None else None
+        row["finish_to_upload_ms"] = uploaded - finished if finished is not None and uploaded is not None else None
+
+    gpu_rows = sorted(gpu_batches.values(), key=lambda row: int(row["batch"]))
+    frame_rows = sorted(frames.values(), key=lambda row: int(row["frame"]))
+    cpu_columns = [
+        "key", "cpu_batch", "queued_ms", "queued_frame", "started_ms", "started_frame",
+        "finished_ms", "finished_frame", "worker_ms", "uploaded_ms", "uploaded_frame",
+        "resident_ms", "resident_frame", "queue_wait_ms", "finish_to_upload_ms",
+    ]
+    gpu_columns = [
+        "batch", "tile_count", "prepared_ms", "prepared_frame", "submitted_ms", "submitted_frame",
+        "completed_ms", "completed_frame", "uploaded_ms", "uploaded_frame", "resident_ms", "resident_frame",
+    ]
+    frame_columns = [
+        "frame", "begin_ms", "end_ms", "inter_draw_gap_ms", "cpu_pump_ms", "gpu_pump_ms",
+        "pre_present_ms", "present_ms", "present_sync", "held",
+    ]
+    cpu_path = prefix.with_name(prefix.name + "_cpu_tiles.csv")
+    gpu_path = prefix.with_name(prefix.name + "_gpu_batches.csv")
+    frame_path = prefix.with_name(prefix.name + "_frames.csv")
+    write_csv(cpu_path, cpu_columns, cpu_rows)
+    write_csv(gpu_path, gpu_columns, gpu_rows)
+    write_csv(frame_path, frame_columns, frame_rows)
+
+    for event in ("queued", "started", "finished", "uploaded", "resident"):
+        count = sum(f"{event}_ms" in row for row in cpu_rows)
+        print(f"cpu_{event}_tiles={count}")
+    print(f"gpu_batches={len(gpu_rows)} gpu_batch_tiles={sum(int(row.get('tile_count', 0)) for row in gpu_rows)} frames={len(frame_rows)}")
+    if cpu_rows:
+        finished_rows = [row for row in cpu_rows if "finished_ms" in row]
+        uploaded_rows = [row for row in cpu_rows if "uploaded_ms" in row]
+        if finished_rows:
+            last_finished = max(finished_rows, key=lambda row: float(row["finished_ms"]))
+            print(f"last_cpu_finished_ms={last_finished['finished_ms']:.3f} key={last_finished['key']}")
+        if uploaded_rows:
+            last_uploaded = max(uploaded_rows, key=lambda row: float(row["uploaded_ms"]))
+            print(f"last_cpu_uploaded_ms={last_uploaded['uploaded_ms']:.3f} key={last_uploaded['key']}")
+        delayed = [row for row in uploaded_rows if row.get("finish_to_upload_ms") is not None]
+        if delayed:
+            max_delay = max(delayed, key=lambda row: float(row["finish_to_upload_ms"]))
+            print(f"max_cpu_finish_to_upload_ms={max_delay['finish_to_upload_ms']:.3f} key={max_delay['key']}")
+    if gpu_rows:
+        for stage in ("prepared", "submitted", "completed", "resident"):
+            present = [row for row in gpu_rows if f"{stage}_ms" in row]
+            if present:
+                last = max(present, key=lambda row: float(row[f"{stage}_ms"]))
+                print(f"last_gpu_batch_{stage}_ms={last[f'{stage}_ms']:.3f} batch={last['batch']} frame={last.get(f'{stage}_frame')}")
+        print(f"cpu_tiles_csv={cpu_path}")
+        print(f"gpu_batches_csv={gpu_path}")
+        print(f"frames_csv={frame_path}")
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())

Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep15Captures.py

@@ -0,0 +1,59 @@
+from pathlib import Path
+
+import numpy as np
+from PIL import Image
+
+
+ROOT = Path(__file__).resolve().parent
+VIEWS = (
+    "chamber",
+    "tunnel",
+    "passage",
+    "far",
+    "passage_floor_closeup",
+    "tunnel_bend",
+    "remote_yaw0",
+    "remote_yaw90",
+    "remote_yaw180",
+    "remote_yaw270",
+)
+
+rows = []
+for view in VIEWS:
+    cpu_path = ROOT / f"step15_cpu_{view}.png"
+    gpu_path = ROOT / f"step15_gpu_{view}.png"
+    cpu = np.asarray(Image.open(cpu_path).convert("RGB"), dtype=np.int16)
+    gpu = np.asarray(Image.open(gpu_path).convert("RGB"), dtype=np.int16)
+    if cpu.shape != gpu.shape:
+        raise RuntimeError(f"dimension mismatch for {view}: {cpu.shape} vs {gpu.shape}")
+    diff = np.abs(cpu - gpu)
+    pixel_max = diff.max(axis=2)
+    rows.append(
+        (
+            view,
+            f"{cpu.shape[1]}x{cpu.shape[0]}",
+            float(diff.mean()),
+            float(np.sqrt(np.square(diff, dtype=np.float64).mean())),
+            float(np.count_nonzero(pixel_max) * 100.0 / pixel_max.size),
+            float(np.count_nonzero(pixel_max > 32) * 100.0 / pixel_max.size),
+            int(pixel_max.max()),
+        )
+    )
+
+lines = [
+    "# Step 15 full-residency capture comparison",
+    "",
+    "All ten CPU/GPU pairs were captured after `[InitialClipmapReady]`. The files were compared numerically as RGB arrays and were not opened in an image viewer.",
+    "",
+    "| View | Size | RGB MAE | RGB RMSE | Changed pixels | Pixels max channel >32 | Max channel |",
+    "|---|---:|---:|---:|---:|---:|---:|",
+]
+for row in rows:
+    view, size, mae, rmse, changed, over32, maximum = row
+    lines.append(
+        f"| {view} | {size} | {mae:.4f} | {rmse:.4f} | {changed:.4f}% | {over32:.4f}% | {maximum} |"
+    )
+
+report = ROOT / "step15_capture_comparison.md"
+report.write_text("\n".join(lines) + "\n", encoding="utf-8")
+print(report.read_text(encoding="utf-8"), end="")

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15CaptureMatrix.ps1

@@ -0,0 +1,80 @@
+[CmdletBinding()]
+param([string]$ExecutableName='VoxelForgeGpuFieldStep15Timeline.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'}
+)
+$EightGB=8L*1024L*1024L*1024L
+foreach($View in $Views)
+{
+    foreach($Mode in @('cpu','gpu'))
+    {
+        $Stem="step15_${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 "STEP15_CAPTURE mode=$Mode view=$($View.Name) affinity=0xFF00 command=`"$ExecutablePath`" $($Arguments -join ' ')"
+        for($Attempt=0;$Attempt-lt2;$Attempt++)
+        {
+            if($Attempt-gt0)
+            {Write-Output '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
+            $PeakPrivate=0L;$OverMemoryLimit=$false
+            try
+            {
+                while(-not$Child.HasExited)
+                {
+                    $Child.Refresh();if($Child.HasExited){break}
+                    $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+                    $PeakPrivate=[Math]::Max($PeakPrivate,$Child.PrivateMemorySize64)
+                    if($PeakPrivate-ge$EightGB)
+                    {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+                    Start-Sleep -Milliseconds 100
+                }
+            }
+            finally
+            {
+                if(-not$Child.HasExited){$Child.Kill();$Child.WaitForExit()}
+            }
+            $Child.WaitForExit();$Child.Refresh();$ChildExitCode=[int]$Child.ExitCode
+            if($ChildExitCode-eq-1073741502-and$Attempt-eq0){continue}
+            if($ChildExitCode-eq-1073741502){throw '0xC0000142 repeated after its single permitted retry.'}
+            if($OverMemoryLimit){throw "Capture viewer stopped at the 8 GiB private-memory ceiling: $Stem"}
+            $LogText=Get-Content -LiteralPath $StdoutPath -Raw
+            if(-not$LogText.Contains('[InitialClipmapReady]')-or-not(Test-Path -LiteralPath $CapturePath -PathType Leaf))
+            {throw "Full-residency capture missing: $Stem; exit=$ChildExitCode"}
+            $Memory=[regex]::Match($LogText,'(?m)^\[GpuFarMemory\] dxgi_local_sampled_peak_mb=([0-9.]+)')
+            $CombinedGB=$PeakPrivate/1GB
+            if($Memory.Success){$CombinedGB=($PeakPrivate+[double]$Memory.Groups[1].Value*1MB)/1GB}
+            if($CombinedGB-ge8.0){throw "Capture exceeded conservative private + DXGI-local 8 GiB ceiling: $Stem ($CombinedGB GiB)"}
+            $Ready=[regex]::Match($LogText,'(?m)^\[InitialClipmapReady\] ([0-9.]+) s')
+            Write-Output "STEP15_CAPTURE_DONE mode=$Mode view=$($View.Name) ready_s=$($Ready.Groups[1].Value) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) conservative_private_plus_dxgi_gb=$([math]::Round($CombinedGB,2)) path=$CapturePath exit=$ChildExitCode"
+            if($ChildExitCode-ne0){throw "Capture viewer returned $ChildExitCode`: $Stem"}
+            break
+        }
+    }
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15CpuReference.ps1

@@ -0,0 +1,63 @@
+[CmdletBinding()]
+param([string]$ExecutableName='VoxelForgeGpuFieldStep15Timeline.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"}
+$Reference=$null
+foreach($Workers in @(1,8))
+{
+    $Name="step15_cpu_reference_w$Workers"
+    $StdoutPath=Join-Path $PSScriptRoot ($Name+'.stdout.log')
+    $StderrPath=Join-Path $PSScriptRoot ($Name+'.stderr.log')
+    $Arguments=@('--stream-benchmark','--shape-world','--force-generate','--stream-hash',
+        '--clip-radius=3',"--workers=$Workers")
+    Write-Output "STEP15_CPU_REFERENCE workers=$Workers affinity=0xFF00 priority=BelowNormal command=`"$ExecutablePath`" $($Arguments -join ' ')"
+    for($Attempt=0;$Attempt-lt2;$Attempt++)
+    {
+        if($Attempt-gt0)
+        {Write-Output '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
+        $PeakPrivate=0L;$OverMemoryLimit=$false
+        try
+        {
+            while(-not$Child.HasExited)
+            {
+                $Child.Refresh();if($Child.HasExited){break}
+                $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+                $PeakPrivate=[Math]::Max($PeakPrivate,$Child.PrivateMemorySize64)
+                if($Child.PrivateMemorySize64-ge(8L*1GB))
+                {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+                Start-Sleep -Milliseconds 100
+            }
+        }
+        finally
+        {
+            if(-not$Child.HasExited){$Child.Kill();$Child.WaitForExit()}
+        }
+        $Child.WaitForExit();$Child.Refresh();$ChildExitCode=[int]$Child.ExitCode
+        if($ChildExitCode-eq-1073741502-and$Attempt-eq0){continue}
+        if($ChildExitCode-eq-1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
+        if($OverMemoryLimit){throw "CPU reference stopped at the 8 GiB private-memory ceiling: W$Workers"}
+        $LogText=Get-Content -LiteralPath $StdoutPath -Raw
+        $Stream=[regex]::Match($LogText,'(?m)^\[StreamFullClipmap\].*generated=2226/2226.*density_hash=(0x[0-9a-fA-F]+) mesh_hash=(0x[0-9a-fA-F]+) payload_hash=(0x[0-9a-fA-F]+) topology_hash=(0x[0-9a-fA-F]+)')
+        $Walkability=[regex]::Match($LogText,'(?m)^\[Walkability\] PASS 5642/5642')
+        if(-not$Stream.Success-or-not$Walkability.Success)
+        {throw "CPU reference failed full-generation/hash/walkability checks for W$Workers; exit=$ChildExitCode"}
+        $Hashes=@($Stream.Groups[1].Value.ToLowerInvariant(),$Stream.Groups[2].Value.ToLowerInvariant(),
+            $Stream.Groups[3].Value.ToLowerInvariant(),$Stream.Groups[4].Value.ToLowerInvariant())
+        if($null-eq$Reference){$Reference=$Hashes}
+        elseif(($Reference-join'|')-ne($Hashes-join'|')){throw 'CPU hashes differ between W1 and W8.'}
+        Write-Output "STEP15_CPU_REFERENCE_RESULT workers=$Workers hashes=$($Hashes -join ',') walkability=5642/5642 peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) exit=$ChildExitCode"
+        if($ChildExitCode-ne0){throw "CPU reference returned $ChildExitCode`: W$Workers"}
+        break
+    }
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15GpuSurvival.ps1

@@ -0,0 +1,59 @@
+[CmdletBinding()]
+param([string]$ExecutableName='VoxelForgeGpuFieldStep15Timeline.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"}
+$Arguments=@('--clip-radius=3','--workers=8','--quit-after-clipmap','--gpu-far-audit','--gpu-far-prewarm')
+$Name='step15_gpu_survival'
+$StdoutPath=Join-Path $PSScriptRoot ($Name+'.stdout.log')
+$StderrPath=Join-Path $PSScriptRoot ($Name+'.stderr.log')
+$EightGB=8L*1024L*1024L*1024L
+Write-Output "STEP15_SURVIVAL_RUN command=`"$ExecutablePath`" $($Arguments -join ' ')"
+for($Attempt=0;$Attempt-lt2;$Attempt++)
+{
+    if($Attempt-gt0)
+    {Write-Output '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
+    $PeakPrivate=0L;$OverMemoryLimit=$false
+    try
+    {
+        while(-not$Child.HasExited)
+        {
+            $Child.Refresh();if($Child.HasExited){break}
+            $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+            $PeakPrivate=[Math]::Max($PeakPrivate,$Child.PrivateMemorySize64)
+            if($Child.PrivateMemorySize64-ge$EightGB)
+            {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+            Start-Sleep -Milliseconds 100
+        }
+    }
+    finally
+    {
+        if(-not$Child.HasExited){$Child.Kill();$Child.WaitForExit()}
+    }
+    $Child.WaitForExit();$Child.Refresh();$ChildExitCode=[int]$Child.ExitCode
+    if($ChildExitCode-eq-1073741502-and$Attempt-eq0){continue}
+    if($ChildExitCode-eq-1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
+    if($OverMemoryLimit){throw 'GPU survival audit stopped at the private-memory ceiling.'}
+    $LogText=Get-Content -LiteralPath $StdoutPath -Raw
+    if(-not$LogText.Contains('[InitialClipmapReady]')){throw "GPU survival audit did not reach full residency; exit=$ChildExitCode"}
+    $Audit=[regex]::Match($LogText,'(?m)^\[GpuFarSurvivalAudit\].*status=(PASS|FAIL)')
+    $Memory=[regex]::Match($LogText,'(?m)^\[GpuFarMemory\] dxgi_local_sampled_peak_mb=([0-9.]+)')
+    if(-not$Audit.Success){throw 'GPU survival audit result marker missing.'}
+    $DxgiMB=if($Memory.Success){[double]$Memory.Groups[1].Value}else{0.0}
+    $CombinedGB=($PeakPrivate+$DxgiMB*1MB)/1GB
+    if($CombinedGB-ge8.0){throw "GPU survival exceeded conservative private + DXGI-local 8 GiB ceiling: $CombinedGB GiB"}
+    Write-Output "STEP15_SURVIVAL_RESULT audit=$($Audit.Groups[1].Value) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) dxgi_local_peak_mb=$DxgiMB conservative_private_plus_dxgi_gb=$([math]::Round($CombinedGB,2)) exit=$ChildExitCode stdout=$StdoutPath stderr=$StderrPath"
+    if($Audit.Groups[1].Value-ne'PASS'){throw 'GPU survival audit failed.'}
+    if($ChildExitCode-ne0){throw "Viewer returned $ChildExitCode"}
+    break
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15Timeline.ps1

@@ -0,0 +1,80 @@
+[CmdletBinding()]
+param(
+    [ValidateSet('Paced','NoVsync','Fast','Race')][string]$Mode='Paced',
+    [string]$ExecutableName='VoxelForgeGpuFieldStep15Timeline.exe'
+)
+
+$ErrorActionPreference='Stop'
+$HostProcess=Get-Process -Id $PID
+$HostProcess.PriorityClass='BelowNormal'
+$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+if([IO.Path]::GetFileName($ExecutableName)-ne$ExecutableName-or[IO.Path]::GetExtension($ExecutableName)-ne'.exe')
+{throw 'ExecutableName must be an executable filename.'}
+$ExecutablePath=Join-Path $StandaloneRoot ('bin\'+$ExecutableName)
+if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "Viewer not found: $ExecutablePath"}
+
+$Runs=switch($Mode)
+{
+    'Paced' {@(@{Mode='gpu';Name='step15_timeline_paced';Extra=@('--gpu-far-no-cold-drain','--gpu-far-timeline')})}
+    'NoVsync' {@(@{Mode='gpu';Name='step15_timeline_no_vsync';Extra=@('--gpu-far-no-cold-drain','--no-vsync','--gpu-far-timeline')})}
+    'Fast' {@(@{Mode='gpu';Name='step15_timeline_fast';Extra=@('--gpu-far-timeline')})}
+    'Race' {@(
+        @{Mode='cpu';Name='step15_race_cpu_A';Extra=@()},
+        @{Mode='gpu';Name='step15_race_gpu_A';Extra=@()},
+        @{Mode='gpu';Name='step15_race_gpu_B';Extra=@()},
+        @{Mode='cpu';Name='step15_race_cpu_B';Extra=@()}
+    )}
+}
+$EightGB=8L*1024L*1024L*1024L
+foreach($Run in $Runs)
+{
+    $Arguments=@('--clip-radius=3','--workers=8','--quit-after-clipmap',"--$($Run.Mode)-far")+$Run.Extra
+    $StdoutPath=Join-Path $PSScriptRoot ($Run.Name+'.stdout.log')
+    $StderrPath=Join-Path $PSScriptRoot ($Run.Name+'.stderr.log')
+    Write-Output "STEP15_RUN mode=$($Run.Mode) name=$($Run.Name) priority=BelowNormal affinity=0xFF00 workers=8 command=`"$ExecutablePath`" $($Arguments -join ' ')"
+    for($Attempt=0;$Attempt-lt2;$Attempt++)
+    {
+        if($Attempt-gt0)
+        {Write-Output '0xC0000142 encountered; waiting 60 seconds before the single permitted retry.';Start-Sleep -Seconds 60}
+        Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
+        $HostRunClock=[Diagnostics.Stopwatch]::StartNew()
+        $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
+            -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
+        $PeakPrivate=0L;$OverMemoryLimit=$false
+        try
+        {
+            while(-not$Child.HasExited)
+            {
+                $Child.Refresh();if($Child.HasExited){break}
+                $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+                $PeakPrivate=[Math]::Max($PeakPrivate,$Child.PrivateMemorySize64)
+                if($Child.PrivateMemorySize64-ge$EightGB)
+                {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+                Start-Sleep -Milliseconds 50
+            }
+        }
+        finally
+        {
+            if(-not$Child.HasExited){$Child.Kill();$Child.WaitForExit()}
+        }
+        $Child.WaitForExit();$Child.Refresh();$HostRunClock.Stop();$HostRunMs=$HostRunClock.Elapsed.TotalMilliseconds
+        $ViewerExitCode=[int]$Child.ExitCode
+        if($ViewerExitCode-eq-1073741502-and$Attempt-eq0){continue}
+        if($ViewerExitCode-eq-1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
+        if($OverMemoryLimit){throw "Viewer stopped at the 8 GiB private-memory ceiling: $($Run.Name)"}
+        $LogText=Get-Content -LiteralPath $StdoutPath -Raw
+        if(-not$LogText.Contains('[InitialClipmapReady]'))
+        {throw "Full clipmap was not ready in $($Run.Name); exit=$ViewerExitCode"}
+        $Ready=[regex]::Match($LogText,'(?m)^\[InitialClipmapReady\] ([0-9.]+) s')
+        $Presented=[regex]::Match($LogText,'(?m)^\[WholeViewPresented\] at_ms=([0-9.]+) present_ms=([0-9.]+).*hr=0x([0-9A-Fa-f]{8})')
+        if(-not$Presented.Success){throw "Whole-view present marker missing in $($Run.Name)"}
+        $Memory=[regex]::Match($LogText,'(?m)^\[GpuFarMemory\] dxgi_local_sampled_peak_mb=([0-9.]+)')
+        $CombinedGB=$PeakPrivate/1GB
+        if($Memory.Success){$CombinedGB=($PeakPrivate+[double]$Memory.Groups[1].Value*1MB)/1GB}
+        if($CombinedGB-ge8.0){throw "Viewer exceeded conservative private + DXGI-local 8 GiB ceiling: $($Run.Name) ($CombinedGB GiB)"}
+        Write-Output "STEP15_RESULT mode=$($Run.Mode) name=$($Run.Name) ready_ms=$([math]::Round([double]$Ready.Groups[1].Value*1000,3)) viewer_whole_view_present_return_ms=$($Presented.Groups[1].Value) present_ms=$($Presented.Groups[2].Value) present_hr=0x$($Presented.Groups[3].Value.ToUpperInvariant()) host_process_exit_ms=$([math]::Round($HostRunMs,3)) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) conservative_private_plus_dxgi_gb=$([math]::Round($CombinedGB,2)) exit=$ViewerExitCode"
+        if($ViewerExitCode-ne0){throw "Viewer returned $ViewerExitCode`: $($Run.Name)"}
+        break
+    }
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/STEP15_REPORT.md

@@ -0,0 +1,49 @@
+# GPU field port — Step 15 report
+
+## Result
+
+Cold-start draining removes most of the post-export delay, and CPU fallback generation now starts alongside GPU input preparation. The interleaved race still shows the hybrid behind the CPU viewer: median first whole view is **1.677 s GPU vs 1.393 s CPU**. The hybrid does not beat the roughly 1.35 s CPU result yet.
+
+## Timeline and diagnosis
+
+The timeline logger is opt-in (`--gpu-far-timeline`). It writes per-frame CSVs, one row per GPU batch, and one row per CPU fallback tile. Times are milliseconds from cold process start; each event also records its frame ordinal. GPU `uploaded` means the compute output became available in its preallocated GPU output buffer; this path has no CPU readback or copy-upload stage.
+
+| Run | First whole view | CPU fallback | GPU batches | Last GPU batch resident |
+|---|---:|---:|---:|---:|
+| Paced control (`--gpu-far-no-cold-drain`) | 2.141 s | 1,144 queued, started, finished, uploaded and resident | 35 / 1,082 tiles | 2.141 s |
+| Paced control + `--no-vsync` | 2.132 s | same workload | 35 / 1,082 tiles | 2.131 s |
+| Cold drain, instrumented | 1.681 s | all 1,144 resident by 1.386 s | 1 / 1,082 tiles | 1.681 s |
+
+In the paced trace, CPU streaming starts at 0.995 s. All 1,144 fallback jobs have finished by 1.205 s, but the last CPU upload/residency event is at 1.949 s. The worst finish-to-upload delay is 755 ms. The upload budget of four batches or 2 ms per frame is the CPU fallback tail. The GPU pump starts around 1.378 s and submits its 35th batch at 2.077 s; that batch completes at 2.140 s and is resident at 2.141 s. The paced schedule spreads its work over 40 frame boundaries.
+
+Disabling vsync only improves the paced control by 9 ms. Startup presentation already uses nonblocking `sync=0`; the cold upload and submission budgets still pace the work. Thus vsync is not the cause of the roughly 0.8 s CPU upload tail.
+
+With cold draining enabled, the instrumented trace has the last CPU fallback tile finished at 1.198 s and resident at 1.386 s. The single GPU batch is prepared at 1.440 s (frame 2), submitted at 1.519 s (frame 2), complete/output-available at 1.670 s (frame 12), and resident at 1.681 s (frame 12). The reported GPU dispatch timestamp total is 197 ms for 1,082 tiles. The residual floor after the input export is therefore the GPU path: command preparation/submission followed by about 200 ms of GPU work and completion/residency polling. In the non-instrumented race, export ends at about 1.33 s and the first whole view arrives 0.32–0.38 s later.
+
+Per-event evidence is in [paced GPU batches](step15_timeline_paced_gpu_batches.csv), [paced CPU tiles](step15_timeline_paced_cpu_tiles.csv), [paced frames](step15_timeline_paced_frames.csv), and the corresponding [cold-drain GPU batches](step15_timeline_fast_gpu_batches.csv), [CPU tiles](step15_timeline_fast_cpu_tiles.csv), and [frames](step15_timeline_fast_frames.csv). The full run logs are [paced](step15_timeline_paced.stdout.log), [no-vsync](step15_timeline_no_vsync.stdout.log), and [cold drain](step15_timeline_fast.stdout.log). Re-run the event summary with `py -3 AnalyzeStep15Timeline.py <log> --prefix <output-prefix>`.
+
+## Changes
+
+- Start the CPU streamer before launching asynchronous GPU input preparation. The CPU fallback jobs can run while GPU inputs are exported and prepared.
+- Before the first complete view only, drain GPU submissions and CPU dirty uploads without the normal frame-sliced caps. Keep the existing four-batch / 2 ms upload budget and GPU submission pacing after the first view is ready. The GPU survival audit retains its bounded submission mode.
+- While waiting for the initial view in cold-drain mode, use a 1 ms wait interval instead of 11 ms.
+- Add `--gpu-far-no-cold-drain` as the paced control and `--gpu-far-timeline` for event/frame telemetry.
+
+The viewer and streamer changes are in `Source/StandaloneViewer.cpp` and `Source/StreamingWorld.h`. The runner and analysis scripts are `RunStep15Timeline.ps1` and `AnalyzeStep15Timeline.py`.
+
+## Interleaved race and remaining floor
+
+The two CPU runs reached the first whole view in 1.422 s and 1.363 s (median **1.393 s**). The GPU runs took 1.647 s and 1.707 s (median **1.677 s**), a **285 ms median gap**. Process-exit medians show the same ordering: 1.628 s CPU vs 1.833 s GPU. The GPU runs ended input export at 1.326–1.331 s, then needed another 0.32–0.38 s to make the full view ready.
+
+GPU shape compilation took 888–897 ms in the race, compared with 585–586 ms in the CPU runs. The logs show asynchronous GPU shader warm-up overlapping shape compilation. This is a likely source of the roughly 300 ms difference, although this race alone does not prove the cause. Moving CPU fallback earlier removed its upload tail as the final gate; it did not make the GPU path faster than the CPU reference. The best measured GPU whole-view time is 1.647 s in the race (1.682 s with detailed event logging).
+
+## Correctness and resource checks
+
+- CPU reference hashes match at one and eight workers: density `0xdb2991c73215af4c`, mesh `0x8236c5edd7fc8e4f`, payload `0x4c48fa75a543d58e`, topology `0xbea22e76b34a4dfc`. Walkability passed **5,642/5,642**. See [one worker](step15_cpu_reference_w1.stdout.log) and [eight workers](step15_cpu_reference_w8.stdout.log).
+- Whole-near and whole-far agreement passed. Both GPU and WARP had zero topology differences and zero protected-exception or thin-feature topology differences in the audited cells. Near: 960 tested tiles, 229,376 protected-exception cells. Far: 1,114 tested tiles, 152 transition tiles skipped, 127,401 protected-exception cells and 353 thin-feature cells. See [near agreement](step15_whole_near_agreement.stdout.log) and [far agreement](step15_whole_far_agreement.stdout.log). This is a sign/topology result, not a claim of bitwise vertex-position parity at every LOD.
+- The GPU survival audit passed: 2.9 GB peak private memory and 2,043 MB peak DXGI local memory (about 4.9 GB combined). Ten full-residency CPU/GPU capture pairs completed; the largest monitored private-plus-DXGI total was 5.06 GB. Captures were compared as numeric RGB arrays and never opened at full size. Per-view values are in [capture comparison](step15_capture_comparison.md).
+- Viewer runs used BelowNormal priority and affinity `0xFF00`; the viewer build used `/MP8`. Runners monitored memory, allowed one retry on `0xC0000142`, and closed the viewers they launched. The step 15 viewer binary is `../../../bin/VoxelForgeGpuFieldStep15Timeline.exe`.
+
+## Follow-up
+
+Direct-GPU LOD3/4 coverage was not changed. The `(6,2,0,lod3)` coverage loss and the resulting 1,144 CPU fallback tiles remain open for a later step. No commit, push, or stash was made.

Standalone/Research/OwnerShapeWorld/gpu_field_port/step15_capture_comparison.md

@@ -0,0 +1,16 @@
+# Step 15 full-residency capture comparison
+
+All ten CPU/GPU pairs were captured after `[InitialClipmapReady]`. The files were compared numerically as RGB arrays and were not opened in an image viewer.
+
+| View | Size | RGB MAE | RGB RMSE | Changed pixels | Pixels max channel >32 | Max channel |
+|---|---:|---:|---:|---:|---:|---:|
+| chamber | 1280x800 | 0.6076 | 1.8655 | 42.6368% | 0.0183% | 124 |
+| tunnel | 1280x800 | 0.6383 | 3.3799 | 31.7335% | 0.2772% | 130 |
+| passage | 1280x800 | 1.6092 | 3.6878 | 50.4982% | 0.1733% | 52 |
+| far | 1280x800 | 0.7909 | 2.1558 | 44.8325% | 0.0198% | 128 |
+| passage_floor_closeup | 1280x800 | 0.9314 | 2.4800 | 43.0219% | 0.0500% | 100 |
+| tunnel_bend | 1280x800 | 0.8129 | 2.1599 | 40.9629% | 0.0708% | 55 |
+| remote_yaw0 | 1280x800 | 0.1994 | 2.6632 | 4.2472% | 0.2415% | 125 |
+| remote_yaw90 | 1280x800 | 0.4083 | 4.3552 | 2.2432% | 0.6948% | 94 |
+| remote_yaw180 | 1280x800 | 0.0484 | 1.3206 | 2.9787% | 0.0341% | 102 |
+| remote_yaw270 | 1280x800 | 0.2355 | 3.1104 | 5.5592% | 0.4098% | 107 |

Standalone/Research/OwnerShapeWorld/gpu_field_port/step15_timeline_fast_frames.csv

@@ -0,0 +1,13 @@
+frame,begin_ms,end_ms,inter_draw_gap_ms,cpu_pump_ms,gpu_pump_ms,pre_present_ms,present_ms,present_sync,held
+1,1100.864,1364.271,0.0,51.311,0.0,0.061,211.608,0,0
+2,1364.285,1528.009,0.014,21.686,136.361,,0.0,,1
+3,1528.013,1543.83,0.004,0.537,0.0,,0.0,,1
+4,1543.835,1559.293,0.005,0.545,0.0,,0.0,,1
+5,1559.297,1575.126,0.005,0.537,0.0,,0.0,,1
+6,1575.129,1590.14,0.003,0.507,0.0,,0.0,,1
+7,1590.146,1606.153,0.006,0.521,0.0,,0.0,,1
+8,1606.158,1621.673,0.005,0.552,0.0,,0.0,,1
+9,1621.678,1637.688,0.005,0.529,0.0,,0.0,,1
+10,1637.69,1653.713,0.002,0.523,0.0,,0.0,,1
+11,1653.718,1669.716,0.005,0.536,0.0,,0.0,,1
+12,1669.721,1682.071,0.005,11.635,0.0,0.167,0.343,1,0

Standalone/Research/OwnerShapeWorld/gpu_field_port/step15_timeline_fast_gpu_batches.csv

@@ -0,0 +1,2 @@
+batch,tile_count,prepared_ms,prepared_frame,submitted_ms,submitted_frame,completed_ms,completed_frame,uploaded_ms,uploaded_frame,resident_ms,resident_frame
+0,1082,1440.022,2,1519.109,2,1669.724,12,1669.724,12,1681.21,12

Standalone/Research/OwnerShapeWorld/gpu_field_port/step15_timeline_paced_frames.csv

@@ -0,0 +1,41 @@
+frame,begin_ms,end_ms,inter_draw_gap_ms,cpu_pump_ms,gpu_pump_ms,pre_present_ms,present_ms,present_sync,held
+1,1069.083,1376.481,0.0,2.081,0.0,0.047,305.114,0,0
+2,1376.488,1409.017,0.007,1.364,19.316,,0.0,,1
+3,1409.023,1440.548,0.006,1.169,18.734,,0.0,,1
+4,1440.553,1472.572,0.005,3.803,8.211,,0.0,,1
+5,1472.575,1504.617,0.003,2.169,3.09,,0.0,,1
+6,1504.623,1535.631,0.006,0.915,4.034,,0.0,,1
+7,1535.637,1567.173,0.006,1.241,3.938,,0.0,,1
+8,1567.179,1583.19,0.006,1.088,3.118,,0.0,,1
+9,1583.194,1615.254,0.004,1.139,4.018,,0.0,,1
+10,1615.26,1631.271,0.005,0.72,3.571,,0.0,,1
+11,1631.275,1662.823,0.005,2.21,3.91,,0.0,,1
+12,1662.828,1694.852,0.005,1.474,5.186,,0.0,,1
+13,1694.861,1710.866,0.009,1.292,0.769,,0.0,,1
+14,1710.87,1742.893,0.004,4.652,0.745,,0.0,,1
+15,1742.898,1758.922,0.005,1.641,0.612,,0.0,,1
+16,1758.927,1774.093,0.005,0.722,0.586,,0.0,,1
+17,1774.098,1789.745,0.005,1.673,0.607,,0.0,,1
+18,1789.749,1805.812,0.004,1.276,0.595,,0.0,,1
+19,1805.815,1821.773,0.003,1.832,0.557,,0.0,,1
+20,1821.778,1837.331,0.006,1.339,0.578,,0.0,,1
+21,1837.333,1853.345,0.003,0.894,0.665,,0.0,,1
+22,1853.351,1869.359,0.005,1.692,0.594,,0.0,,1
+23,1869.362,1885.372,0.004,1.783,0.522,,0.0,,1
+24,1885.46,1901.387,0.088,1.441,0.619,,0.0,,1
+25,1901.389,1916.91,0.003,0.697,0.638,,0.0,,1
+26,1916.916,1932.443,0.006,2.677,0.607,,0.0,,1
+27,1932.448,1948.454,0.005,1.129,0.654,,0.0,,1
+28,1948.456,1963.974,0.003,0.85,0.665,,0.0,,1
+29,1963.979,1979.984,0.006,0.597,0.782,,0.0,,1
+30,1979.987,1996.0,0.003,0.511,0.774,,0.0,,1
+31,1996.003,2012.013,0.004,0.555,0.751,,0.0,,1
+32,2012.015,2028.032,0.002,0.211,0.809,,0.0,,1
+33,2028.035,2043.55,0.002,0.589,0.698,,0.0,,1
+34,2043.552,2059.84,0.002,0.714,0.636,,0.0,,1
+35,2059.846,2075.579,0.006,1.253,0.504,,0.0,,1
+36,2075.67,2091.594,0.09,0.486,0.532,,0.0,,1
+37,2091.601,2107.883,0.007,0.13,0.0,,0.0,,1
+38,2107.89,2123.633,0.006,0.19,0.0,,0.0,,1
+39,2123.64,2139.648,0.006,0.593,0.0,,0.0,,1
+40,2139.653,2141.213,0.005,0.967,0.0,0.164,0.214,1,0

Standalone/Research/OwnerShapeWorld/gpu_field_port/step15_timeline_paced_gpu_batches.csv

@@ -0,0 +1,36 @@
+batch,tile_count,prepared_ms,prepared_frame,submitted_ms,submitted_frame,completed_ms,completed_frame,uploaded_ms,uploaded_frame,resident_ms,resident_frame
+0,1,1388.247,2,1397.152,2,1428.933,3,1428.933,3,1428.944,3
+1,32,1413.788,3,1428.779,3,1452.575,4,1452.575,4,1452.962,4
+2,32,1446.663,4,1452.387,4,1477.841,5,1477.841,5,1478.167,5
+3,32,1476.904,5,1477.716,5,1509.58,6,1509.58,6,1509.904,6
+4,32,1508.634,6,1509.455,6,1540.823,7,1540.823,7,1541.164,7
+5,32,1539.89,7,1540.698,7,1571.391,8,1571.391,8,1571.745,8
+6,32,1570.524,8,1571.348,8,1588.361,9,1588.361,9,1588.831,9
+7,32,1587.36,9,1588.213,9,1619.568,10,1619.568,10,1619.906,10
+8,32,1618.58,10,1619.437,10,1637.409,11,1637.409,11,1637.768,11
+9,32,1636.46,11,1637.279,11,1669.497,12,1669.497,12,1669.875,12
+10,32,1668.525,12,1669.335,12,1712.223,14,1712.223,14,1715.47,14
+11,32,1696.536,13,1696.773,13,1790.566,18,1790.566,18,1790.965,18
+12,32,1716.002,14,1716.131,14,1806.677,19,1806.677,19,1807.244,19
+13,32,1744.93,15,1745.036,15,1806.908,19,1806.908,19,1807.551,19
+14,32,1760.01,16,1760.119,16,1807.254,19,1807.254,19,1807.607,19
+15,32,1776.158,17,1776.262,17,1822.736,20,1822.736,20,1823.076,20
+16,32,1791.398,18,1791.506,18,1870.466,23,1870.466,23,1870.791,23
+17,32,1807.989,19,1808.168,19,1870.794,23,1870.794,23,1871.105,23
+18,32,1823.473,20,1823.582,20,1886.083,24,1886.083,24,1886.458,24
+19,32,1838.669,21,1838.779,21,1886.461,24,1886.461,24,1886.87,24
+20,32,1855.403,22,1855.522,22,1918.875,26,1918.875,26,1919.237,26
+21,32,1871.524,23,1871.632,23,1919.24,26,1919.24,26,1919.564,26
+22,32,1887.293,24,1887.406,24,1933.218,27,1933.218,27,1933.548,27
+23,32,1902.493,25,1902.611,25,1964.111,29,1964.111,29,1964.532,29
+24,32,1919.974,26,1920.09,26,1980.078,30,1980.078,30,1980.454,30
+25,32,1933.992,27,1934.117,27,1996.101,31,1996.101,31,1996.514,31
+26,32,1949.667,28,1949.805,28,2028.126,33,2028.126,33,2028.542,33
+27,32,1965.075,29,1965.229,29,2043.643,34,2043.643,34,2044.176,34
+28,32,1980.984,30,1981.137,30,2059.958,35,2059.958,35,2060.398,35
+29,32,1996.98,31,1997.17,31,2060.401,35,2060.401,35,2060.745,35
+30,32,2012.661,32,2012.887,32,2060.748,35,2060.748,35,2061.077,35
+31,32,2029.032,33,2029.187,33,2075.774,36,2075.774,36,2076.128,36
+32,32,2044.651,34,2044.782,34,2123.753,39,2123.753,39,2124.177,39
+33,32,2061.439,35,2061.568,35,2139.765,40,2139.765,40,2140.204,40
+34,25,2076.466,36,2076.657,36,2140.207,40,2140.207,40,2140.534,40

Standalone/Source/StandaloneViewer.cpp

@@ -1317,6 +1317,10 @@ namespace
         uint8 ActiveSkirtFaceMask=0x3f;
         uint64 Samples=0;
         double WorkerSeconds=0.0;
+        std::chrono::steady_clock::time_point CpuQueuedAt{};
+        std::chrono::steady_clock::time_point CpuStartedAt{};
+        std::chrono::steady_clock::time_point CpuFinishedAt{};
+        bool bColdTimelineResidentReported=false;
     };
     struct FStreamingGpuBatch
     {
@@ -1365,17 +1369,25 @@ namespace
         uint8 VisibleSkirtFaceMask=0;
         bool bSkirtDrawArgsAudited=false;
         double GpuDispatchMs=0.0;
+        std::chrono::steady_clock::time_point PreparedAt{};
+        std::chrono::steady_clock::time_point SubmittedAt{};
+        std::chrono::steady_clock::time_point CompletedAt{};
         bool bGpuComplete=false;
         bool bTimingResolved=false;
+        bool bColdTimelineResidentReported=false;
     };
     struct FGpuFarBatch
     {
         ID3D11Query* CompletionEvent=nullptr;
         VoxelForgeGpuFarViewer::GpuTimingQueries TimingQueries{};
         uint32 TileCount=0;
+        uint32 ResidentCount=0;
+        std::vector<FVoxelTileKey> Keys;
         uint64 TimestampFrequency=0;
         double GpuDispatchMs=0.0;
+        std::chrono::steady_clock::time_point PreparedAt{};
         std::chrono::steady_clock::time_point SubmittedAt{};
+        std::chrono::steady_clock::time_point CompletedAt{};
         bool bGpuComplete=false,bTimingResolved=false;
     };
 
@@ -1487,6 +1499,7 @@ namespace
         bool bShapeCompare=false;
         bool bFarShapeStagePending=false;
         bool bGpuFarEnabled=false;
+        bool bDisableGpuColdStartDrain=false;
         bool bGpuFarAuditEnabled=false;
         bool bGpuFarShaderPrewarm=false,bDisableGpuWarmupPresentHold=false;
         bool bGpuFarWarmupResourcesOnly=false;
@@ -1495,6 +1508,10 @@ namespace
         bool bGpuFarWarmupOnly=false;
         std::string GpuWarmupEntryFilter;
         std::thread GpuShaderWarmupThread;
+        std::thread GpuFarPreparationThread;
+        std::atomic<bool> bGpuFarPreparationFinished{false};
+        bool bGpuFarPreparationStarted=false,bGpuFarInputsFinalized=false;
+        bool bGpuWarmupCompleted=false,bGpuFarTimelineEnabled=false;
         bool bGpuShaderWarmupQueueSucceeded=false;
         std::string GpuShaderWarmupQueueError;
         bool bGpuFarAuditReported=false;
@@ -1529,6 +1546,8 @@ namespace
         std::vector<uint32> GpuFarFreeOutputSlots;
         std::unordered_map<FVoxelTileKey,VoxelForgeCore::FarFieldTileInput,FStandaloneTileKeyHash> GpuFarPreparedInputs;
         std::unordered_map<FVoxelTileKey,std::string,FStandaloneTileKeyHash> GpuFarPreparationErrors;
+        std::unordered_map<FVoxelTileKey,std::chrono::steady_clock::time_point,FStandaloneTileKeyHash> GpuFarPreparedAt;
+        std::vector<FVoxelTileKey> GpuFarCandidateTiles,GpuFarCpuFallbackTiles,GpuFarNeedPreparation;
         double GpuFarPreparationWallMs=0.0,GpuFarPreparationWorkerMs=0.0;
         uint64 GpuFarPumpCalls=0,GpuFarCompletionPolls=0;
         double GpuFarFirstPumpMs=-1.0,GpuFarLastPumpMs=-1.0;
@@ -1568,6 +1587,7 @@ namespace
         std::string CapturePath;
         std::string MeshCachePath;
         std::chrono::steady_clock::time_point ColdProcessStart{};
+        std::atomic<uint64> ColdTimelineFrameOrdinal{0};
         std::chrono::steady_clock::time_point GenerationStart{};
         bool bShowLevel[MaxLevels+1]={true,true,true,true,true,true,true,true,true};
         bool Keys[256]{};
@@ -1595,6 +1615,7 @@ namespace
         double BatchUploadBudgetMilliseconds=2.0;
 
         ~FApp(){
+            if(GpuFarPreparationThread.joinable())GpuFarPreparationThread.join();
             if(GpuShaderWarmupThread.joinable())GpuShaderWarmupThread.join();
             StopGeneration(true);ReleaseMeshes();
             if(SolidRasterizer)SolidRasterizer->Release();if(WireRasterizer)WireRasterizer->Release();
@@ -1619,6 +1640,15 @@ namespace
             Tile={};
         }
 
+        double ColdTimelineMilliseconds(std::chrono::steady_clock::time_point At)const
+        {
+            return At.time_since_epoch().count()==0?-1.0:
+                std::chrono::duration<double,std::milli>(At-ColdProcessStart).count();
+        }
+
+        uint64 ColdTimelineFrame()const
+        {return ColdTimelineFrameOrdinal.load(std::memory_order_relaxed);}
+
         static void ReleaseGpuFarBatch(FGpuFarBatch& Batch)
         {
             if(Batch.CompletionEvent)Batch.CompletionEvent->Release();
@@ -1654,7 +1684,11 @@ namespace
             for(FGpuFarBatch& Batch:GpuFarBatches)ReleaseGpuFarBatch(Batch);
             GpuFarBatches.clear();GpuFarQueue.clear();GpuFarQueued.clear();GpuFarEligibleTiles.clear();
             GpuFarFreeOutputSlots.clear();GpuFarNextOutputSlot=0;
-            GpuFarPreparedInputs.clear();GpuFarPreparationErrors.clear();
+            GpuFarPreparedInputs.clear();GpuFarPreparationErrors.clear();GpuFarPreparedAt.clear();
+            GpuFarCandidateTiles.clear();GpuFarCpuFallbackTiles.clear();GpuFarNeedPreparation.clear();
+            bGpuFarPreparationStarted=false;bGpuFarInputsFinalized=false;
+            bGpuFarPreparationFinished.store(false,std::memory_order_relaxed);
+            bGpuWarmupCompleted=false;
             GpuFarPreparationWallMs=0.0;GpuFarPreparationWorkerMs=0.0;
             GpuFarPumpCalls=0;GpuFarCompletionPolls=0;
             GpuFarFirstPumpMs=GpuFarLastPumpMs=GpuFarFirstSubmitMs=GpuFarLastSubmitMs=-1.0;
@@ -1683,6 +1717,8 @@ namespace
             (void)bCancel;
             if(Streamer)Streamer->Stop();
             Streamer.reset();
+            if(GpuFarPreparationThread.joinable())GpuFarPreparationThread.join();
+            bGpuFarPreparationStarted=false;
             if(FullShapeBuildThread.joinable())FullShapeBuildThread.join();
             PendingFullShapeWorld.reset();
             bFullShapeBuildFinished.store(false,std::memory_order_relaxed);
@@ -2259,14 +2295,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 (!bShapeWorld&&MeshCache.IsLoaded())?&MeshCache:nullptr,
                 bFarShapeStagePending?World->FullResClipLevels-1:-1,
                 false,false,false,false,ShapeWorld.get());
+            if(bGpuFarTimelineEnabled)
+                Streamer->EnableColdTimelineTrace(ColdProcessStart,&ColdTimelineFrameOrdinal);
             if(bGpuFarEnabled&&ShapeWorld&&ShapeWorld->hasFarWorld())
             {
                 ScheduleGpuFarTiles(DesiredTiles);
                 if(bGpuFarAuditEnabled)
                 {
-                    std::vector<FVoxelTileKey> AuditKeys;AuditKeys.reserve(GpuFarEligibleTiles.size());
+                    std::vector<FVoxelTileKey> AuditKeys;AuditKeys.reserve(GpuFarCandidateTiles.size());
                     for(const FVoxelTileKey& Key:DesiredTiles)
-                        if(Key.Level>=6&&GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key);
+                        if(Key.Level>=6&&std::find(GpuFarCandidateTiles.begin(),GpuFarCandidateTiles.end(),Key)
+                            !=GpuFarCandidateTiles.end())AuditKeys.push_back(Key);
                     GpuFarAudit=std::make_unique<FFarSurvivalAudit>(*World,*ShapeWorld,AuditKeys,false);
                 }
             }
@@ -2293,6 +2332,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             std::printf("Streaming seed %d%s%s: %d desired tiles, %d full-detail tiles, %d workers; upload budget %d batches / %.1f ms per frame.\n",
                 Seed,TunnelMode,bShapeWorld?" [shape intervals]":"",Update.DesiredTiles,
                 Update.FullResTiles,Streamer->GetWorkerCount(),MaxBatchUploadsPerFrame,BatchUploadBudgetMilliseconds);
+            StartGpuFarPreparation();
             if(bShapePipelineActive)
             {
                 std::string ShapeError;
@@ -2316,6 +2356,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         void StartFarShapeGeneration()
         {
             if(!bFarShapeStagePending||!ShapeWorld||ShapeWorld->hasFarWorld()||!Streamer)return;
+            if(GpuFarPreparationThread.joinable())GpuFarPreparationThread.join();
             ShapeWorld->buildFarWorld();
             bFarShapeStagePending=false;
             Streamer->Stop();
@@ -2324,14 +2365,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             Streamer.reset();
             Streamer=std::make_unique<FStandaloneStreamingGenerator>(*World,WorkerCount,false,nullptr,-1,
                 false,false,false,false,ShapeWorld.get());
+            if(bGpuFarTimelineEnabled)
+                Streamer->EnableColdTimelineTrace(ColdProcessStart,&ColdTimelineFrameOrdinal);
             if(bGpuFarEnabled)
             {
                 ScheduleGpuFarTiles(DesiredTiles);
                 if(bGpuFarAuditEnabled&&World)
                 {
-                    std::vector<FVoxelTileKey> AuditKeys;AuditKeys.reserve(GpuFarEligibleTiles.size());
+                    std::vector<FVoxelTileKey> AuditKeys;AuditKeys.reserve(GpuFarCandidateTiles.size());
                     for(const FVoxelTileKey& Key:DesiredTiles)
-                        if(Key.Level>=6&&GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key);
+                        if(Key.Level>=6&&std::find(GpuFarCandidateTiles.begin(),GpuFarCandidateTiles.end(),Key)
+                            !=GpuFarCandidateTiles.end())AuditKeys.push_back(Key);
                     GpuFarAudit=std::make_unique<FFarSurvivalAudit>(*World,*ShapeWorld,AuditKeys,false);
                 }
             }
@@ -2344,6 +2388,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             bGenerationRunning=true;
             std::printf("[FarStreamStart] desired=%d retained=%d new_jobs=%d workers=%d\n",
                 Update.DesiredTiles,Update.RetainedResidentTiles,Update.NewJobs,WorkerCount);
+            StartGpuFarPreparation();
             std::fflush(stdout);
         }
 
@@ -2359,6 +2404,46 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             }
         }
 
+        void ReportGpuTileCompletion(const FVoxelTileKey& Key,FGpuFarTile& Tile,
+                                     uint32 BatchIndex=UINT32_MAX)
+        {
+            if(Tile.CompletedAt.time_since_epoch().count()!=0)return;
+            Tile.CompletedAt=std::chrono::steady_clock::now();
+            if(!bGpuFarTimelineEnabled)return;
+            std::printf("[GpuTileTimeline] event=completed key=(%d,%d,%d,lod%d) batch=%u at_ms=%.3f frame=%llu submitted_ms=%.3f\n",
+                Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,BatchIndex,
+                ColdTimelineMilliseconds(Tile.CompletedAt),
+                static_cast<unsigned long long>(ColdTimelineFrame()),
+                ColdTimelineMilliseconds(Tile.SubmittedAt));
+            std::printf("[GpuTileTimeline] event=uploaded key=(%d,%d,%d,lod%d) batch=%u at_ms=%.3f frame=%llu method=direct_gpu_compute_output copy_upload=none\n",
+                Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,BatchIndex,
+                ColdTimelineMilliseconds(Tile.CompletedAt),
+                static_cast<unsigned long long>(ColdTimelineFrame()));
+        }
+
+        void ReportGpuTileResident(const FVoxelTileKey& Key,FGpuFarTile& Tile)
+        {
+            if(Tile.bColdTimelineResidentReported)return;
+            Tile.bColdTimelineResidentReported=true;
+            const auto ResidentAt=std::chrono::steady_clock::now();
+            if(bGpuFarTimelineEnabled)
+                std::printf("[GpuTileTimeline] event=resident key=(%d,%d,%d,lod%d) batch=%u at_ms=%.3f frame=%llu submitted_ms=%.3f completed_ms=%.3f\n",
+                    Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,Tile.BatchIndex,
+                    ColdTimelineMilliseconds(ResidentAt),
+                    static_cast<unsigned long long>(ColdTimelineFrame()),
+                    ColdTimelineMilliseconds(Tile.SubmittedAt),ColdTimelineMilliseconds(Tile.CompletedAt));
+            if(Tile.BatchIndex!=UINT32_MAX&&Tile.BatchIndex<GpuFarBatches.size())
+            {
+                FGpuFarBatch& Batch=GpuFarBatches[Tile.BatchIndex];
+                ++Batch.ResidentCount;
+                if(bGpuFarTimelineEnabled&&Batch.ResidentCount==Batch.TileCount)
+                    std::printf("[GpuBatchTimeline] event=resident batch=%u at_ms=%.3f frame=%llu tile_count=%u submitted_ms=%.3f completed_ms=%.3f\n",
+                        Tile.BatchIndex,ColdTimelineMilliseconds(ResidentAt),
+                        static_cast<unsigned long long>(ColdTimelineFrame()),Batch.TileCount,
+                        ColdTimelineMilliseconds(Batch.SubmittedAt),ColdTimelineMilliseconds(Batch.CompletedAt));
+            }
+        }
+
         bool IsRenderedTileReady(const FVoxelTileKey& Key)
         {
             if(bGpuFarEnabled&&ShapeWorld&&ShapeWorld->hasFarWorld()
@@ -2375,10 +2460,30 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     {
                         Batch.bGpuComplete=true;
                         const auto PollAt=std::chrono::steady_clock::now();
+                        Batch.CompletedAt=PollAt;
                         const double PollMs=std::chrono::duration<double,std::milli>(PollAt-ColdProcessStart).count();
                         if(GpuFarFirstCompletionPollMs<0.0)GpuFarFirstCompletionPollMs=PollMs;
                         GpuFarLastCompletionPollMs=PollMs;++GpuFarCompletionPolls;
                         Batch.CompletionEvent->Release();Batch.CompletionEvent=nullptr;
+                        if(bGpuFarTimelineEnabled)
+                        {
+                            std::printf("[GpuBatchTimeline] event=completed batch=%u at_ms=%.3f frame=%llu tile_count=%u submitted_ms=%.3f\n",
+                                Tile.BatchIndex,ColdTimelineMilliseconds(PollAt),
+                                static_cast<unsigned long long>(ColdTimelineFrame()),Batch.TileCount,
+                                ColdTimelineMilliseconds(Batch.SubmittedAt));
+                            std::printf("[GpuBatchTimeline] event=uploaded batch=%u at_ms=%.3f frame=%llu method=direct_gpu_compute_output copy_upload=none\n",
+                                Tile.BatchIndex,ColdTimelineMilliseconds(PollAt),
+                                static_cast<unsigned long long>(ColdTimelineFrame()));
+                        }
+                        for(const FVoxelTileKey& BatchKey:Batch.Keys)
+                        {
+                            auto BatchTileIt=GpuFarTiles.find(BatchKey);
+                            if(BatchTileIt!=GpuFarTiles.end())
+                            {
+                                BatchTileIt->second.bGpuComplete=true;
+                                ReportGpuTileCompletion(BatchKey,BatchTileIt->second,Tile.BatchIndex);
+                            }
+                        }
                     }
                     if(Batch.bGpuComplete&&!Batch.bTimingResolved)
                     {
@@ -2402,7 +2507,11 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                             GpuFarGpuDispatchMaxMs=std::max(GpuFarGpuDispatchMaxMs,TileDispatchMs);
                         }
                     }
-                    if(Batch.bGpuComplete)RecordGpuCrossingTileReady(Key);
+                    if(Batch.bGpuComplete)
+                    {
+                        RecordGpuCrossingTileReady(Key);
+                        ReportGpuTileResident(Key,Tile);
+                    }
                     return Batch.bGpuComplete;
                 }
                 if(!Tile.bGpuComplete&&Tile.CompletionEvent&&Context
@@ -2410,6 +2519,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 {
                     Tile.bGpuComplete=true;
                     Tile.CompletionEvent->Release();Tile.CompletionEvent=nullptr;
+                    ReportGpuTileCompletion(Key,Tile);
                 }
                 if(Tile.bGpuComplete&&!Tile.bTimingResolved)
                 {
@@ -2425,7 +2535,11 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                         Tile.bTimingResolved=true;
                     }
                 }
-                if(Tile.bGpuComplete)RecordGpuCrossingTileReady(Key);
+                if(Tile.bGpuComplete)
+                {
+                    RecordGpuCrossingTileReady(Key);
+                    ReportGpuTileResident(Key,Tile);
+                }
                 return Tile.bGpuComplete;
             }
             return ResidentTiles.contains(Key);
@@ -2436,9 +2550,11 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             if(NeedPreparation.empty())return;
             const auto PrepareStart=std::chrono::steady_clock::now();
             std::printf("[ColdTimeline] stage=gpu_input_export begin_ms=%.3f tiles=%zu\n",
-                std::chrono::duration<double,std::milli>(PrepareStart-ColdProcessStart).count(),NeedPreparation.size());
+                ColdTimelineMilliseconds(PrepareStart),NeedPreparation.size());
             std::vector<std::optional<VoxelForgeCore::FarFieldTileInput>> Inputs(NeedPreparation.size());
             std::vector<std::string> Errors(NeedPreparation.size());
+            std::vector<std::chrono::steady_clock::time_point> FinishedAt(NeedPreparation.size());
+            std::vector<uint64> FinishedFrame(NeedPreparation.size());
             const std::size_t ThreadCount=std::min<std::size_t>(8,NeedPreparation.size());
             std::atomic<std::size_t> Next{0};
             std::vector<double> WorkerElapsed(ThreadCount,0.0);
@@ -2475,6 +2591,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                             }
                         }
                         catch(const std::exception& Error) { Errors[Index]=Error.what(); }
+                        FinishedAt[Index]=std::chrono::steady_clock::now();
+                        FinishedFrame[Index]=ColdTimelineFrame();
                     }
                     WorkerElapsed[Worker]=std::chrono::duration<double,std::milli>(
                         std::chrono::steady_clock::now()-WorkerStart).count();
@@ -2488,7 +2606,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             std::size_t Failed=0;
             for(std::size_t Index=0;Index<NeedPreparation.size();++Index)
             {
-                if(Inputs[Index])GpuFarPreparedInputs.emplace(NeedPreparation[Index],std::move(*Inputs[Index]));
+                if(Inputs[Index])
+                {
+                    GpuFarPreparedInputs.emplace(NeedPreparation[Index],std::move(*Inputs[Index]));
+                    GpuFarPreparedAt.emplace(NeedPreparation[Index],FinishedAt[Index]);
+                    const FVoxelTileKey& Key=NeedPreparation[Index];
+                    if(bGpuFarTimelineEnabled)
+                        std::printf("[GpuTileTimeline] event=prepared key=(%d,%d,%d,lod%d) at_ms=%.3f frame=%llu\n",
+                            Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,
+                            ColdTimelineMilliseconds(FinishedAt[Index]),
+                            static_cast<unsigned long long>(FinishedFrame[Index]));
+                }
                 else
                 {
                     ++Failed;
@@ -2499,7 +2627,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             std::printf("[GpuFarInputPrepare] tiles=%zu workers=%zu wall_ms=%.3f worker_sum_ms=%.3f failed=%zu\n",
                 NeedPreparation.size(),ThreadCount,PreparationMs,WorkerSumMs,Failed);
             std::printf("[ColdTimeline] stage=gpu_input_export end_ms=%.3f duration_ms=%.3f worker_sum_ms=%.3f\n",
-                std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count(),
+                ColdTimelineMilliseconds(std::chrono::steady_clock::now()),
                 PreparationMs,WorkerSumMs);
             std::fflush(stdout);
         }
@@ -2507,27 +2635,11 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         void ScheduleGpuFarTiles(const std::vector<FVoxelTileKey>& Keys)
         {
             GpuFarQueue.clear();GpuFarQueued.clear();GpuFarEligibleTiles.clear();
-            GpuFarPreparedInputs.clear();GpuFarPreparationErrors.clear();
+            GpuFarPreparedInputs.clear();GpuFarPreparationErrors.clear();GpuFarPreparedAt.clear();
+            GpuFarCandidateTiles.clear();GpuFarCpuFallbackTiles.clear();GpuFarNeedPreparation.clear();
+            bGpuFarPreparationStarted=false;bGpuFarInputsFinalized=false;
+            bGpuFarPreparationFinished.store(false,std::memory_order_relaxed);
             if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld())return;
-            if(bGpuFarShaderPrewarm)
-            {
-                if(GpuShaderWarmupThread.joinable())GpuShaderWarmupThread.join();
-                if(!bGpuShaderWarmupQueueSucceeded)
-                {
-                    std::fprintf(stderr,"[GpuShaderWarmupError] queue: %s\n",
-                        GpuShaderWarmupQueueError.empty()?"warmup queue failed":GpuShaderWarmupQueueError.c_str());
-                    bGpuFarEnabled=false;
-                    return;
-                }
-                std::string WarmupError;
-                if(!VoxelForgeGpuFarViewer::CompleteFarShaderWarmup(Device,Context,
-                    GpuFarShaderPath,WarmupError))
-                {
-                    std::fprintf(stderr,"[GpuShaderWarmupError] %s\n",WarmupError.c_str());
-                    bGpuFarEnabled=false;
-                    return;
-                }
-            }
             std::vector<FVoxelTileKey> CpuTiles;CpuTiles.reserve(Keys.size());
             std::vector<FVoxelTileKey> Eligible;Eligible.reserve(Keys.size());
             uint64 SkirtTransitionFallbackTiles=0;
@@ -2606,49 +2718,102 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 std::printf("[CaptureOutputSlotProbe] slot=%d pool_slots=%u\n",
                     CaptureOutputSlot,GpuFarOutputCapacitySlots);
             }
-            if(!Eligible.empty())
+            GpuFarCandidateTiles=std::move(Eligible);
+            GpuFarCpuFallbackTiles=std::move(CpuTiles);
+            GpuFarNeedPreparation.reserve(GpuFarCandidateTiles.size());
+            for(const FVoxelTileKey& Key:GpuFarCandidateTiles)
+                if(!GpuFarTiles.contains(Key))GpuFarNeedPreparation.push_back(Key);
+            if(Streamer)Streamer->SetDesiredTileFilter(GpuFarCpuFallbackTiles);
+            const std::size_t NearGpuTiles=static_cast<std::size_t>(std::count_if(
+                GpuFarCandidateTiles.begin(),GpuFarCandidateTiles.end(),[&](const FVoxelTileKey& Key)
+                    {return Key.Level<World->FullResClipLevels;}));
+            std::printf("[GpuFieldEligibility] candidate_gpu_tiles=%zu near_gpu_tiles=%zu cpu_fallback_tiles=%zu\n",
+                GpuFarCandidateTiles.size(),NearGpuTiles,GpuFarCpuFallbackTiles.size());
+            std::printf("[GpuSkirtSeamFallback] near_tiles=%llu reason=fine_visible_lod_skirts gpu_far_skirts=active\n",
+                static_cast<unsigned long long>(SkirtTransitionFallbackTiles));
+        }
+
+        void AddGpuFarCpuFallbackTiles(const std::vector<FVoxelTileKey>& Additional)
         {
-                std::string PoolError;const auto ReserveStart=std::chrono::steady_clock::now();
+            if(Additional.empty()||!Streamer)return;
+            std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> Seen(
+                GpuFarCpuFallbackTiles.begin(),GpuFarCpuFallbackTiles.end());
+            std::size_t Added=0;
+            for(const FVoxelTileKey& Key:Additional)
+                if(Seen.insert(Key).second){GpuFarCpuFallbackTiles.push_back(Key);++Added;}
+            if(Added==0)return;
+            Streamer->SetDesiredTileFilter(GpuFarCpuFallbackTiles);
+            const FVector CameraVoxels(Position.X*4.0,Position.Y*4.0,Position.Z*4.0);
+            std::vector<FVoxelTileKey> ResidentKeys;ResidentKeys.reserve(ResidentTiles.size());
+            for(const auto& Pair:ResidentTiles)ResidentKeys.push_back(Pair.first);
+            const FStandaloneStreamingUpdate Update=Streamer->UpdateDesired(
+                CameraVoxels,ResidentKeys,CrossingIndex);
+            FCrossingWork& Work=CrossingWork[CrossingIndex];
+            Work.SubmittedJobs+=Update.NewJobs;Work.PendingJobs+=Update.NewJobs;
+            std::printf("[GpuFarPrepFallbackAdded] tiles=%zu queued_jobs=%d cpu_fallback_total=%zu\n",
+                Added,Update.NewJobs,GpuFarCpuFallbackTiles.size());
+        }
+
+        void StartGpuFarPreparation()
+        {
+            if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld()||bGpuFarPreparationStarted)return;
+            bGpuFarPreparationStarted=true;
+            if(GpuFarCandidateTiles.empty())
+            {
+                bGpuFarInputsFinalized=true;
+                bGpuFarPreparationFinished.store(true,std::memory_order_relaxed);
+                return;
+            }
+            const auto ReserveStart=std::chrono::steady_clock::now();
             std::printf("[ColdTimeline] stage=gpu_output_pool begin_ms=%.3f slots=%u\n",
-                    std::chrono::duration<double,std::milli>(ReserveStart-ColdProcessStart).count(),
+                ColdTimelineMilliseconds(ReserveStart),GpuFarOutputCapacitySlots);
-                    GpuFarOutputCapacitySlots);
+            std::string PoolError;
             if(!VoxelForgeGpuFarViewer::ReserveFarOutputPool(Device,Context,GpuFarShaderPath,
                 GpuFarOutputCapacitySlots,static_cast<uint32>(World->CoarseTileCells),PoolError))
             {
                 std::fprintf(stderr,"[GpuFarPoolError] %s\n",PoolError.c_str());
                 bGpuFarEnabled=false;GpuFarOutputCapacitySlots=0;
-                    CpuTiles=Keys;Eligible.clear();GpuFarTiles.clear();
+                for(auto& Pair:GpuFarTiles)ReleaseGpuFarTile(Pair.second);
+                GpuFarTiles.clear();GpuFarCandidateTiles.clear();GpuFarQueue.clear();GpuFarQueued.clear();
+                GpuFarNeedPreparation.clear();GpuFarEligibleTiles.clear();
+                AddGpuFarCpuFallbackTiles(DesiredTiles);
+                bGpuFarInputsFinalized=true;
+                bGpuFarPreparationFinished.store(true,std::memory_order_relaxed);
+                return;
             }
             GpuFarPoolReserveMs+=std::chrono::duration<double,std::milli>(
                 std::chrono::steady_clock::now()-ReserveStart).count();
             std::printf("[ColdTimeline] stage=gpu_output_pool end_ms=%.3f duration_ms=%.3f\n",
-                    std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count(),
+                ColdTimelineMilliseconds(std::chrono::steady_clock::now()),GpuFarPoolReserveMs);
-                    GpuFarPoolReserveMs);
             SampleGpuFarMemory();
-            }
+            if(GpuFarNeedPreparation.empty())
-            std::vector<FVoxelTileKey> NeedPreparation;NeedPreparation.reserve(Eligible.size());
-            for(const FVoxelTileKey& Key:Eligible)
-                if(!GpuFarTiles.contains(Key))NeedPreparation.push_back(Key);
-            if(!NeedPreparation.empty())
-                PrepareGpuFarInputs(std::move(NeedPreparation));
-            std::vector<FVoxelTileKey> PreparedEligible;PreparedEligible.reserve(Eligible.size());
-            for(const FVoxelTileKey& Key:Eligible)
             {
-                if(!GpuFarTiles.contains(Key)&&!GpuFarPreparedInputs.contains(Key))
+                bGpuFarPreparationFinished.store(true,std::memory_order_relaxed);
+                return;
+            }
+            bGpuFarPreparationFinished.store(false,std::memory_order_relaxed);
+            const std::vector<FVoxelTileKey> NeedPreparation=GpuFarNeedPreparation;
+            GpuFarPreparationThread=std::thread([this,NeedPreparation]()
             {
-                    CpuTiles.push_back(Key);
+                try { PrepareGpuFarInputs(NeedPreparation); }
-                    const auto Error=GpuFarPreparationErrors.find(Key);
+                catch(const std::exception& Error)
-                    std::fprintf(stderr,"[GpuFieldInputFallback] key=(%d,%d,%d,lod%d) reason=%s\n",
+                {
-                        Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,
+                    for(const FVoxelTileKey& Key:NeedPreparation)
-                        Error==GpuFarPreparationErrors.end()?"missing prepared input":Error->second.c_str());
+                        GpuFarPreparationErrors.emplace(Key,Error.what());
-                    continue;
+                    std::fprintf(stderr,"[GpuFarInputPrepareError] %s\n",Error.what());
                 }
-                GpuFarEligibleTiles.insert(Key);
+                catch(...)
-                if(!GpuFarTiles.contains(Key))
+                {
-                {GpuFarQueue.push_back(Key);GpuFarQueued.insert(Key);}
+                    for(const FVoxelTileKey& Key:NeedPreparation)
-                PreparedEligible.push_back(Key);
+                        GpuFarPreparationErrors.emplace(Key,"unknown input preparation failure");
+                    std::fprintf(stderr,"[GpuFarInputPrepareError] unknown failure\n");
                 }
-            Eligible.swap(PreparedEligible);
+                bGpuFarPreparationFinished.store(true,std::memory_order_release);
+            });
+        }
+
+        void SortGpuFarQueue()
+        {
             std::sort(GpuFarQueue.begin(),GpuFarQueue.end(),[&](const FVoxelTileKey& A,const FVoxelTileKey& B)
             {
                 if(A.Level!=B.Level)return A.Level<B.Level;
@@ -2662,20 +2827,79 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 return AC<BC;
             });
             GpuFarQueued.clear();for(const FVoxelTileKey& Key:GpuFarQueue)GpuFarQueued.insert(Key);
-            if(Streamer)Streamer->SetDesiredTileFilter(CpuTiles);
+        }
+
+        bool FinalizeGpuFarPreparation()
+        {
+            if(bGpuFarInputsFinalized)return true;
+            if(!bGpuFarPreparationStarted
+                ||!bGpuFarPreparationFinished.load(std::memory_order_acquire))return false;
+            if(GpuFarPreparationThread.joinable())GpuFarPreparationThread.join();
+            if(bGpuFarShaderPrewarm&&!bGpuWarmupCompleted)
+            {
+                if(GpuShaderWarmupThread.joinable())GpuShaderWarmupThread.join();
+                if(!bGpuShaderWarmupQueueSucceeded)
+                {
+                    std::fprintf(stderr,"[GpuShaderWarmupError] queue: %s\n",
+                        GpuShaderWarmupQueueError.empty()?"warmup queue failed":GpuShaderWarmupQueueError.c_str());
+                    bGpuFarEnabled=false;
+                }
+                else
+                {
+                    std::string WarmupError;
+                    if(!VoxelForgeGpuFarViewer::CompleteFarShaderWarmup(Device,Context,
+                        GpuFarShaderPath,WarmupError))
+                    {
+                        std::fprintf(stderr,"[GpuShaderWarmupError] %s\n",WarmupError.c_str());
+                        bGpuFarEnabled=false;
+                    }
+                    else bGpuWarmupCompleted=true;
+                }
+            }
+            if(!bGpuFarEnabled)
+            {
+                GpuFarCandidateTiles.clear();GpuFarQueue.clear();GpuFarQueued.clear();
+                GpuFarEligibleTiles.clear();
+                for(auto& Pair:GpuFarTiles)ReleaseGpuFarTile(Pair.second);
+                GpuFarTiles.clear();
+                AddGpuFarCpuFallbackTiles(DesiredTiles);
+                bGpuFarInputsFinalized=true;
+                return false;
+            }
+            std::vector<FVoxelTileKey> Failed;
+            for(const FVoxelTileKey& Key:GpuFarCandidateTiles)
+            {
+                if(GpuFarTiles.contains(Key))
+                {GpuFarEligibleTiles.insert(Key);continue;}
+                if(!GpuFarPreparedInputs.contains(Key))
+                {
+                    Failed.push_back(Key);
+                    const auto Error=GpuFarPreparationErrors.find(Key);
+                    std::fprintf(stderr,"[GpuFieldInputFallback] key=(%d,%d,%d,lod%d) reason=%s\n",
+                        Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,
+                        Error==GpuFarPreparationErrors.end()?"missing prepared input":Error->second.c_str());
+                    continue;
+                }
+                GpuFarEligibleTiles.insert(Key);
+                GpuFarQueue.push_back(Key);
+            }
+            SortGpuFarQueue();
+            AddGpuFarCpuFallbackTiles(Failed);
+            bGpuFarInputsFinalized=true;
             const std::size_t NearGpuTiles=static_cast<std::size_t>(std::count_if(
                 GpuFarEligibleTiles.begin(),GpuFarEligibleTiles.end(),[&](const FVoxelTileKey& Key)
                     {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());
+                GpuFarEligibleTiles.size(),NearGpuTiles,GpuFarCpuFallbackTiles.size());
-            std::printf("[GpuSkirtSeamFallback] near_tiles=%llu reason=fine_visible_lod_skirts gpu_far_skirts=active\n",
+            return true;
-                static_cast<unsigned long long>(SkirtTransitionFallbackTiles));
         }
 
         void PumpGpuFarGeneration()
         {
             if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld()||!Device||!Context
-                ||GpuFarQueue.empty())return;
+                ||!FinalizeGpuFarPreparation()||GpuFarQueue.empty())return;
+            const bool bColdDrainBeforeWholeView=!bInitialClipmapReported&&!bGpuFarAuditEnabled
+                &&!bDisableGpuColdStartDrain;
             const auto PumpStart=std::chrono::steady_clock::now();
             const double PumpMs=std::chrono::duration<double,std::milli>(PumpStart-ColdProcessStart).count();
             if(GpuFarFirstPumpMs<0.0)GpuFarFirstPumpMs=PumpMs;
@@ -2697,7 +2921,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             }
             const auto Start=std::chrono::steady_clock::now();
             int32 Built=0;
-            std::vector<FVoxelTileKey> BuiltKeys;BuiltKeys.reserve(32);
+            std::vector<FVoxelTileKey> BuiltKeys;
+            BuiltKeys.reserve(bColdDrainBeforeWholeView?GpuFarQueue.size():32);
             do
             {
                 const FVoxelTileKey Key=GpuFarQueue.front();GpuFarQueue.pop_front();GpuFarQueued.erase(Key);
@@ -2762,6 +2987,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 Tile.VertexCapacityBytes=static_cast<uint64>(
                     VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile)*32u;
                 Tile.VisibleSkirtFaceMask=PreparedSkirtFaceMask;
+                const auto PreparedAtIt=GpuFarPreparedAt.find(Key);
+                if(PreparedAtIt!=GpuFarPreparedAt.end())Tile.PreparedAt=PreparedAtIt->second;
+                Tile.SubmittedAt=std::chrono::steady_clock::now();
                 Tile.GpuDispatchMs=GpuDispatchMs;
                 GpuFarEnqueueWallSeconds+=EnqueueMs/1000.0;
                 if(GpuDispatchMs>0.0)
@@ -2773,6 +3001,12 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 ++GpuFarTilesSubmitted;++Built;
                 GpuFarTiles.emplace(Key,Tile);
                 BuiltKeys.push_back(Key);
+                if(!bDeferredBatch&&bGpuFarTimelineEnabled)
+                    std::printf("[GpuTileTimeline] event=submitted key=(%d,%d,%d,lod%d) batch=none at_ms=%.3f frame=%llu prepared_ms=%.3f path=direct_dispatch\n",
+                        Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,
+                        ColdTimelineMilliseconds(Tile.SubmittedAt),
+                        static_cast<unsigned long long>(ColdTimelineFrame()),
+                        ColdTimelineMilliseconds(Tile.PreparedAt));
                 if(!bGpuFarFirstTileReported)
                 {
                     bGpuFarFirstTileReported=true;
@@ -2782,13 +3016,21 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     std::fflush(stdout);
                 }
             }
-            while(Built<32&&!GpuFarQueue.empty()
+            while(!GpuFarQueue.empty()&&(bColdDrainBeforeWholeView
-                &&std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-Start).count()<4.0);
+                ||(Built<32&&std::chrono::duration<double,std::milli>(
+                    std::chrono::steady_clock::now()-Start).count()<4.0)));
             if(Built>0)
             {
                 if(bDeferredBatch)
                 {
                     FGpuFarBatch Batch;Batch.TileCount=static_cast<uint32>(Built);
+                    Batch.PreparedAt=std::chrono::steady_clock::now();
+                    Batch.Keys=BuiltKeys;
+                    const uint32 BatchIndex=static_cast<uint32>(GpuFarBatches.size());
+                    if(bGpuFarTimelineEnabled)
+                        std::printf("[GpuBatchTimeline] event=prepared batch=%u at_ms=%.3f frame=%llu tile_count=%u path=direct_gpu_output\n",
+                            BatchIndex,ColdTimelineMilliseconds(Batch.PreparedAt),
+                            static_cast<unsigned long long>(ColdTimelineFrame()),Batch.TileCount);
                     std::string BatchError;
                     if(!VoxelForgeGpuFarViewer::SubmitFarTileBatch(Device,Context,GpuFarShaderPath,
                         &Batch.CompletionEvent,&Batch.TimingQueries,BatchError))
@@ -2803,13 +3045,24 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                             Batch.SubmittedAt-ColdProcessStart).count();
                         if(GpuFarFirstSubmitMs<0.0)GpuFarFirstSubmitMs=SubmitMs;
                         GpuFarLastSubmitMs=SubmitMs;
-                        const uint32 BatchIndex=static_cast<uint32>(GpuFarBatches.size());
                         GpuFarBatches.push_back(Batch);
                         for(const FVoxelTileKey& Key:BuiltKeys)
                         {
                             auto It=GpuFarTiles.find(Key);
-                            if(It!=GpuFarTiles.end())It->second.BatchIndex=BatchIndex;
+                            if(It!=GpuFarTiles.end())
+                            {
+                                It->second.BatchIndex=BatchIndex;
+                                It->second.SubmittedAt=Batch.SubmittedAt;
+                                if(bGpuFarTimelineEnabled)std::printf("[GpuTileTimeline] event=submitted key=(%d,%d,%d,lod%d) batch=%u at_ms=%.3f frame=%llu prepared_ms=%.3f\n",
+                                    Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,BatchIndex,
+                                    ColdTimelineMilliseconds(Batch.SubmittedAt),
+                                    static_cast<unsigned long long>(ColdTimelineFrame()),
+                                    ColdTimelineMilliseconds(It->second.PreparedAt));
                             }
+                        }
+                        if(bGpuFarTimelineEnabled)std::printf("[GpuBatchTimeline] event=submitted batch=%u at_ms=%.3f frame=%llu tile_count=%u\n",
+                            BatchIndex,ColdTimelineMilliseconds(Batch.SubmittedAt),
+                            static_cast<unsigned long long>(ColdTimelineFrame()),Batch.TileCount);
                         ++GpuFarSubmissionBatches;
                         GpuFarMaxTilesPerSubmissionBatch=std::max<uint32>(GpuFarMaxTilesPerSubmissionBatch,
                             static_cast<uint32>(Built));
@@ -3219,14 +3472,46 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             };
             std::sort(Dirty.begin(),Dirty.end(),[&](const auto& A,const auto& B)
             {const double DA=Distance(A),DB=Distance(B);return DA==DB?A.Level<B.Level:DA<DB;});
+            const bool bDrainColdStart=bGpuFarEnabled&&!bInitialClipmapReported
+                &&!bDisableGpuColdStartDrain;
             const auto Start=std::chrono::steady_clock::now();int32 Uploaded=0;
             for(const FStandaloneBatchKey& Key:Dirty)
             {
-                if(Uploaded>=MaxBatchUploadsPerFrame)break;
+                if(!bDrainColdStart&&Uploaded>=MaxBatchUploadsPerFrame)break;
-                if(Uploaded>0&&BatchUploadBudgetMilliseconds>0.0
+                if(!bDrainColdStart&&Uploaded>0&&BatchUploadBudgetMilliseconds>0.0
                     &&std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-Start).count()>=BatchUploadBudgetMilliseconds)break;
                 if(!DirtyBatchSet.contains(Key))continue;
-                if(RebuildBatch(Key)){DirtyBatchSet.erase(Key);++Uploaded;}
+                if(RebuildBatch(Key))
+                {
+                    DirtyBatchSet.erase(Key);++Uploaded;
+                    if(bGpuFarTimelineEnabled)
+                    {
+                        const auto UploadAt=std::chrono::steady_clock::now();
+                        const bool bHasUploadBuffer=StreamingBatches.contains(Key);
+                        std::printf("[CpuUploadBatchTimeline] event=uploaded batch=(lod%d,%d,%d,%d) at_ms=%.3f frame=%llu has_buffer=%s\n",
+                            Key.Level,Key.X,Key.Y,Key.Z,ColdTimelineMilliseconds(UploadAt),
+                            static_cast<unsigned long long>(ColdTimelineFrame()),
+                            bHasUploadBuffer?"yes":"no_empty_mesh");
+                        const auto MembersIt=BatchMembers.find(Key);
+                        if(MembersIt!=BatchMembers.end())for(const FVoxelTileKey& TileKey:MembersIt->second)
+                        {
+                            auto ResidentIt=ResidentTiles.find(TileKey);
+                            if(ResidentIt==ResidentTiles.end()||ResidentIt->second.bColdTimelineResidentReported)continue;
+                            FResidentTileMesh& Resident=ResidentIt->second;
+                            Resident.bColdTimelineResidentReported=true;
+                            std::printf("[CpuTileTimeline] event=uploaded key=(%d,%d,%d,lod%d) at_ms=%.3f frame=%llu batch=(lod%d,%d,%d,%d) queued_ms=%.3f started_ms=%.3f finished_ms=%.3f method=%s\n",
+                                TileKey.Coord.X,TileKey.Coord.Y,TileKey.Coord.Z,TileKey.Level,
+                                ColdTimelineMilliseconds(UploadAt),static_cast<unsigned long long>(ColdTimelineFrame()),
+                                Key.Level,Key.X,Key.Y,Key.Z,ColdTimelineMilliseconds(Resident.CpuQueuedAt),
+                                ColdTimelineMilliseconds(Resident.CpuStartedAt),ColdTimelineMilliseconds(Resident.CpuFinishedAt),
+                                bHasUploadBuffer?"cpu_batch_vertex_index":"empty_surface_no_gpu_copy");
+                            std::printf("[CpuTileTimeline] event=resident key=(%d,%d,%d,lod%d) at_ms=%.3f frame=%llu batch=(lod%d,%d,%d,%d)\n",
+                                TileKey.Coord.X,TileKey.Coord.Y,TileKey.Coord.Z,TileKey.Level,
+                                ColdTimelineMilliseconds(UploadAt),static_cast<unsigned long long>(ColdTimelineFrame()),
+                                Key.Level,Key.X,Key.Y,Key.Z);
+                        }
+                    }
+                }
             }
         }
 
@@ -3346,6 +3631,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 Resident.TriangleSkirtFaceMasks=std::move(Tile.TriangleSkirtFaceMasks);
                 Resident.ActiveSkirtFaceMask=ComputeVisibleSkirtFaceMask(Tile.Key,DesiredTileSet);
                 Resident.Samples=Tile.DensitySamples;Resident.WorkerSeconds=Tile.WorkerSeconds;
+                Resident.CpuQueuedAt=Tile.QueuedAt;Resident.CpuStartedAt=Tile.StartedAt;
+                Resident.CpuFinishedAt=Tile.CompletedAt;
                 ResidentTiles.emplace(Tile.Key,std::move(Resident));
                 const FStandaloneBatchKey Batch=BatchForTile(Tile.Key);
                 BatchMembers[Batch].push_back(Tile.Key);MarkBatchDirty(Batch);
@@ -3534,7 +3821,14 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
 
         void Draw()
         {
+            const uint64 FrameOrdinal=ColdTimelineFrameOrdinal.fetch_add(1,std::memory_order_relaxed)+1;
             const auto DrawStart=std::chrono::steady_clock::now();
+            const bool bTraceColdFrame=bGpuFarTimelineEnabled&&!bInitialClipmapReported;
+            if(bTraceColdFrame)
+                std::printf("[ColdFrame] stage=begin frame=%llu at_ms=%.3f inter_draw_gap_ms=%.3f\n",
+                    static_cast<unsigned long long>(FrameOrdinal),ColdTimelineMilliseconds(DrawStart),
+                    LastDrawEndAt.time_since_epoch().count()!=0
+                        ?std::chrono::duration<double,std::milli>(DrawStart-LastDrawEndAt).count():0.0);
             const double InterDrawGapMs=LastDrawEndAt.time_since_epoch().count()!=0
                 ?std::chrono::duration<double,std::milli>(DrawStart-LastDrawEndAt).count():0.0;
             PumpGeneration();
@@ -3706,8 +4000,14 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 ++GpuWarmupHeldPresents;
                 // Keep pumping CPU and GPU work at the display cadence without
                 // calling the driver present path that stalls on its first GPU frame.
-                std::this_thread::sleep_for(std::chrono::milliseconds(11));
+                std::this_thread::sleep_for(std::chrono::milliseconds(
+                    bGpuFarEnabled&&!bDisableGpuColdStartDrain?1:11));
                 LastDrawEndAt=std::chrono::steady_clock::now();
+                if(bTraceColdFrame)
+                    std::printf("[ColdFrame] stage=end frame=%llu at_ms=%.3f cpu_pump_ms=%.3f gpu_pump_ms=%.3f held=1 present_ms=0.000\n",
+                        static_cast<unsigned long long>(FrameOrdinal),ColdTimelineMilliseconds(LastDrawEndAt),
+                        std::chrono::duration<double,std::milli>(CpuPumpEnd-DrawStart).count(),
+                        std::chrono::duration<double,std::milli>(GpuPumpEnd-CpuPumpEnd).count());
                 return;
             }
             const auto BeforePresent=std::chrono::steady_clock::now();
@@ -3786,6 +4086,12 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     InterDrawGapMs,CpuPumpMs,GpuPumpMs,
                     std::chrono::duration<double,std::milli>(BeforePresent-GpuPumpEnd).count(),
                     PresentSyncInterval,PresentCallMs,TitleMs);
+            if(bTraceColdFrame)
+                std::printf("[ColdFrame] stage=end frame=%llu at_ms=%.3f cpu_pump_ms=%.3f gpu_pump_ms=%.3f pre_present_ms=%.3f present_ms=%.3f present_sync=%u present_hr=0x%08lx held=0\n",
+                    static_cast<unsigned long long>(FrameOrdinal),ColdTimelineMilliseconds(DrawEnd),
+                    CpuPumpMs,GpuPumpMs,
+                    std::chrono::duration<double,std::milli>(BeforePresent-GpuPumpEnd).count(),
+                    PresentCallMs,PresentSyncInterval,static_cast<unsigned long>(PresentResult));
             LastDrawEndAt=DrawEnd;
         }
     };
@@ -5632,6 +5938,8 @@ int main(int Argc,char** Argv)
         else if(Arg=="--affinity-all-cpu"){}
         else if(Arg=="--gpu-far-no-present-hold")
             App.bDisableGpuWarmupPresentHold=true;
+        else if(Arg=="--gpu-far-no-cold-drain")App.bDisableGpuColdStartDrain=true;
+        else if(Arg=="--gpu-far-timeline")App.bGpuFarTimelineEnabled=true;
         else if(Arg=="--gpu-far-no-prewarm")
             App.bGpuFarShaderPrewarm=false;
         else if(Arg=="--cpu-far")App.bGpuFarEnabled=false,App.bShapeWorld=true,App.bForceGenerate=true;

Standalone/Source/StreamingWorld.h

@@ -6,6 +6,7 @@
 #include <condition_variable>
 #include <algorithm>
 #include <array>
+#include <atomic>
 #include <bit>
 #include <chrono>
 #include <cstdint>
@@ -41,6 +42,8 @@ struct FStandaloneCpuTile
     uint32 SkirtTriangleCount=0;
     uint64 DensitySamples=0;
     double WorkerSeconds=0.0;
+    std::chrono::steady_clock::time_point QueuedAt{};
+    std::chrono::steady_clock::time_point StartedAt{};
     std::chrono::steady_clock::time_point CompletedAt{};
     uint64 DensityHash=14695981039346656037ull;
     bool bHasDensityHash=false;
@@ -340,6 +343,8 @@ struct FStandaloneTileJob
     int32 BandZMax=MAX_int32;
     uint64 RequestedCrossing=0;
     size_t PriorityRank=0;
+    std::chrono::steady_clock::time_point QueuedAt{};
+    std::chrono::steady_clock::time_point StartedAt{};
 };
 
 struct FStandaloneTileJobHeapOrder
@@ -421,6 +426,14 @@ public:
     FStandaloneStreamingGenerator& operator=(const FStandaloneStreamingGenerator&)=delete;
     ~FStandaloneStreamingGenerator(){Stop();}
 
+    void EnableColdTimelineTrace(std::chrono::steady_clock::time_point Origin,
+                                 const std::atomic<uint64>* FrameOrdinal)
+    {
+        TimelineOrigin=Origin;
+        TimelineFrameOrdinal=FrameOrdinal;
+        bColdTimelineTrace=true;
+    }
+
     FStandaloneStreamingUpdate Start(const FVector& CameraPositionVoxels,
                                      const std::vector<FVoxelTileKey>& ResidentKeys={},
                                      uint64 CrossingIndex=0)
@@ -540,6 +553,7 @@ private:
         auto QueueEnd=std::remove_if(Queue.begin(),Queue.end(),[&](const FStandaloneTileJob& Job)
         {
             if(Desired.contains(Job.Key))return false;
+            if(bColdTimelineTrace)TraceTileEvent("cancelled",Job.Key,FClock::now(),Job.QueuedAt);
             Outstanding.erase(Job.Key);++Stats.CancelledQueuedJobs;
             ++CancelledByCrossing[Job.RequestedCrossing];return true;
         });
@@ -568,7 +582,12 @@ private:
         for(const FVoxelTileKey& Key:NewDesired.Tiles)
         {
             if(Resident.contains(Key)||Outstanding.contains(Key))continue;
-            Queue.push_back({Key,BandMin,BandMax,CrossingIndex,PriorityRanks.at(Key)});
+            FStandaloneTileJob Job;
+            Job.Key=Key;Job.BandZMin=BandMin;Job.BandZMax=BandMax;
+            Job.RequestedCrossing=CrossingIndex;Job.PriorityRank=PriorityRanks.at(Key);
+            Job.QueuedAt=FClock::now();
+            Queue.push_back(Job);
+            if(bColdTimelineTrace)TraceTileEvent("queued",Key,Job.QueuedAt);
             Outstanding.insert(Key);++Stats.NewJobs;
         }
         CurrentDesired=std::move(NewDesired);
@@ -578,8 +597,9 @@ private:
     }
     FStandaloneCpuTile Generate(const FStandaloneTileJob& Job)
     {
-        const auto StartTime=FClock::now();
+        const auto StartTime=Job.StartedAt;
         FStandaloneCpuTile Out;Out.Key=Job.Key;Out.RequestedCrossing=Job.RequestedCrossing;
+        Out.QueuedAt=Job.QueuedAt;Out.StartedAt=Job.StartedAt;
         if(MeshCache)
         {
             if(const FStandaloneCpuTile* Cached=MeshCache->Find(Job.Key))
@@ -718,9 +738,12 @@ private:
                 Wake.wait(Lock,[&](){return bStop||!Queue.empty();});
                 if(bStop)return;
                 std::pop_heap(Queue.begin(),Queue.end(),FStandaloneTileJobHeapOrder{});
-                Job=Queue.back();Queue.pop_back();
+                Job=Queue.back();Queue.pop_back();Job.StartedAt=FClock::now();
             }
+            if(bColdTimelineTrace)TraceTileEvent("started",Job.Key,Job.StartedAt,Job.QueuedAt,Job.StartedAt);
             FStandaloneCpuTile Tile=Generate(Job);
+            if(bColdTimelineTrace)TraceTileEvent("finished",Tile.Key,Tile.CompletedAt,
+                Tile.QueuedAt,Tile.StartedAt,Tile.CompletedAt,Tile.WorkerSeconds*1000.0);
             {
                 std::lock_guard<std::mutex> Lock(Mutex);
                 if(bStop)return;
@@ -730,6 +753,24 @@ private:
     }
 
     using FClock=std::chrono::steady_clock;
+    double TimelineMilliseconds(FClock::time_point At)const
+    {
+        return At.time_since_epoch().count()==0?-1.0:
+            std::chrono::duration<double,std::milli>(At-TimelineOrigin).count();
+    }
+    uint64 TimelineFrame()const
+    {
+        return TimelineFrameOrdinal?TimelineFrameOrdinal->load(std::memory_order_relaxed):0;
+    }
+    void TraceTileEvent(const char* Event,const FVoxelTileKey& Key,FClock::time_point At,
+        FClock::time_point Queued={},FClock::time_point Started={},FClock::time_point Finished={},
+        double WorkerMs=-1.0)const
+    {
+        std::printf("[CpuTileTimeline] event=%s key=(%d,%d,%d,lod%d) at_ms=%.3f frame=%llu queued_ms=%.3f started_ms=%.3f finished_ms=%.3f worker_ms=%.3f\n",
+            Event,Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,TimelineMilliseconds(At),
+            static_cast<unsigned long long>(TimelineFrame()),TimelineMilliseconds(Queued),
+            TimelineMilliseconds(Started),TimelineMilliseconds(Finished),WorkerMs);
+    }
     FStandaloneOwnerWorld* World=nullptr;
     int32 WorkerCount=8;
     bool bHashDensity=false;
@@ -753,4 +794,7 @@ private:
     bool bUseTileFilter=false;
     FStandaloneDesiredTileSet CurrentDesired;
     int32 FullResTileCount=0;
+    bool bColdTimelineTrace=false;
+    FClock::time_point TimelineOrigin{};
+    const std::atomic<uint64>* TimelineFrameOrdinal=nullptr;
 };