diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep15Captures.py b/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep15Captures.py
index 8b21919..e8607d5 100644
--- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep15Captures.py
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep15Captures.py
@@ -21,6 +21,15 @@ VIEWS = (
     "remote_yaw180",
     "remote_yaw270",
 )
+ROCK = np.array((96, 86, 68))
+BACKGROUND = np.array((14, 19, 27))
+
+
+def mask_bounds(mask):
+    ys, xs = np.nonzero(mask)
+    if not len(xs):
+        return "none"
+    return f"({xs.min()},{ys.min()})-({xs.max()},{ys.max()})"
 
 rows = []
 for view in VIEWS:
@@ -32,6 +41,10 @@ for view in VIEWS:
         raise RuntimeError(f"dimension mismatch for {view}: {cpu.shape} vs {gpu.shape}")
     diff = np.abs(cpu - gpu)
     pixel_max = diff.max(axis=2)
+    cpu_rock = np.all(cpu == ROCK, axis=2)
+    gpu_rock = np.all(gpu == ROCK, axis=2)
+    cpu_to_gpu_background = cpu_rock & np.all(gpu == BACKGROUND, axis=2)
+    gpu_to_cpu_background = gpu_rock & np.all(cpu == BACKGROUND, axis=2)
     rows.append(
         (
             view,
@@ -41,6 +54,9 @@ for view in VIEWS:
             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()),
+            int(np.count_nonzero(cpu_to_gpu_background)),
+            int(np.count_nonzero(gpu_to_cpu_background)),
+            mask_bounds(cpu_to_gpu_background),
         )
     )
 
@@ -49,13 +65,15 @@ lines = [
     "",
     "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 |",
-    "|---|---:|---:|---:|---:|---:|---:|",
+    "`CPU rock to GPU background` isolates the prior missing-coverage signature. Captures were loaded as RGB arrays only; no full-size image was opened in a viewer.",
+    "",
+    "| View | Size | RGB MAE | RGB RMSE | Changed pixels | Pixels max channel >32 | Max channel | CPU rock to GPU background | GPU rock to CPU background | Missing-pixel bounds |",
+    "|---|---:|---:|---:|---:|---:|---:|---:|---:|---|",
 ]
 for row in rows:
-    view, size, mae, rmse, changed, over32, maximum = row
+    view, size, mae, rmse, changed, over32, maximum, missing, extra, bounds = row
     lines.append(
-        f"| {view} | {size} | {mae:.4f} | {rmse:.4f} | {changed:.4f}% | {over32:.4f}% | {maximum} |"
+        f"| {view} | {size} | {mae:.4f} | {rmse:.4f} | {changed:.4f}% | {over32:.4f}% | {maximum} | {missing} | {extra} | {bounds} |"
     )
 
 report = ROOT / f"{PREFIX}_capture_comparison.md"
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep17TileControls.py b/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep17TileControls.py
new file mode 100644
index 0000000..643720a
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep17TileControls.py
@@ -0,0 +1,85 @@
+import argparse
+from pathlib import Path
+
+import numpy as np
+from PIL import Image
+
+
+ROOT = Path(__file__).resolve().parent
+DEFAULT_CPU_PATH = "step17_tilecontrol_cpu.png"
+DEFAULT_CONTROLS = (
+    ("GPU indirect draw", "step17_tilecontrol_gpu.png"),
+    ("GPU direct draw", "step17_directdraw_gpu_direct_draw.png"),
+    ("GPU readback copy (cm)", "step17_readbackcopy_gpu_readback_copy.png"),
+    ("GPU readback copy (m)", "step17_metercopy_gpu_meter_copy.png"),
+    ("CPU mesh in GPU draw loop", "step17_cpumeshcopy_gpu_cpu_mesh_copy.png"),
+    ("GPU depth test disabled", "step17_tilecontrol_gpu_depth_disabled.png"),
+    ("GPU skirts disabled", "step17_tilecontrol_gpu_skirts_off.png"),
+    ("GPU connectors disabled", "step17_tilecontrol_gpu_connectors_off.png"),
+    ("GPU skirt-depth fix", "step17_r19_fixed_gpu.png"),
+    (
+        "GPU skirts and connectors disabled",
+        "step17_tilecontrol_gpu_skirts_connectors_off.png",
+    ),
+)
+
+parser = argparse.ArgumentParser()
+parser.add_argument("--cpu", default=DEFAULT_CPU_PATH)
+parser.add_argument("--control", action="append", default=[], metavar="LABEL=PNG")
+parser.add_argument("--report", default="step17_tilecontrol_comparison.md")
+parser.add_argument("--include-existing-controls", action="store_true")
+args = parser.parse_args()
+
+if Path(args.cpu).name != args.cpu or Path(args.report).name != args.report:
+    raise RuntimeError("CPU image and report must be filenames within this directory")
+controls = list(DEFAULT_CONTROLS) if args.include_existing_controls else []
+for item in args.control:
+    label, separator, filename = item.partition("=")
+    if not separator or not label or Path(filename).name != filename:
+        raise RuntimeError(f"invalid --control value: {item!r}")
+    controls.append((label, filename))
+
+CPU_PATH = ROOT / args.cpu
+
+cpu = np.asarray(Image.open(CPU_PATH).convert("RGB"), dtype=np.int16)
+cpu_rock = np.all(cpu == (96, 86, 68), axis=2)
+background = np.array((14, 19, 27))
+rows = []
+for label, filename in controls:
+    path = ROOT / filename
+    gpu = np.asarray(Image.open(path).convert("RGB"), dtype=np.int16)
+    if cpu.shape != gpu.shape:
+        raise RuntimeError(f"dimension mismatch for {label}: {cpu.shape} vs {gpu.shape}")
+    diff = np.abs(cpu - gpu)
+    pixel_max = diff.max(axis=2)
+    missing_cpu_rock = cpu_rock & np.all(gpu == background, axis=2)
+    rows.append(
+        (
+            label,
+            f"{gpu.shape[1]}x{gpu.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),
+            int(np.count_nonzero(missing_cpu_rock)),
+            int(pixel_max.max()),
+        )
+    )
+
+lines = [
+    "# Step 17 LOD3 tile control comparison",
+    "",
+    "The `(6,2,0,lod3)` target was the only drawn tile, with the full desired clipmap retained. Captures were loaded as RGB arrays and compared numerically; full-size images were not opened in a viewer.",
+    "",
+    "`CPU rock to GPU background` counts pixels where the CPU target is rock `(96,86,68)` and the control is background `(14,19,27)`, the prior missing-coverage signature.",
+    "",
+    "| Control | Size | RGB MAE | RGB RMSE | Changed pixels | CPU rock to GPU background | Max channel |",
+    "|---|---:|---:|---:|---:|---:|---:|",
+]
+for label, size, mae, rmse, changed, missing, maximum in rows:
+    lines.append(
+        f"| {label} | {size} | {mae:.4f} | {rmse:.4f} | {changed:.4f}% | {missing} | {maximum} |"
+    )
+
+report = ROOT / args.report
+report.write_text("\n".join(lines) + "\n", encoding="utf-8")
+print(report.read_text(encoding="utf-8"), end="")
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompileStep14Shader.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompileStep14Shader.ps1
new file mode 100644
index 0000000..f6d8727
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompileStep14Shader.ps1
@@ -0,0 +1,29 @@
+[CmdletBinding()]
+param([string]$Entry='ApplyNearFieldDetailCS')
+
+$ErrorActionPreference='Stop'
+$BuildHostProcess=Get-Process -Id $PID
+$BuildHostProcess.PriorityClass='BelowNormal'
+$BuildHostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$ShaderSource=Join-Path $StandaloneRoot 'Research\OwnerShapeWorld\gpu_first_port\VoxelForgeGpuMesher.hlsl'
+$OutputPath=Join-Path $PSScriptRoot ('precompiled\'+$Entry+'.cso')
+if(-not(Test-Path -LiteralPath $ShaderSource -PathType Leaf)){throw "Shader source missing: $ShaderSource"}
+if($Entry-notin@('ApplyNearFieldDetailCS')){throw "Unsupported Step 14 shader entry: $Entry"}
+$FxCRoot='C:\Program Files (x86)\Windows Kits\10\bin'
+$FxC=Get-ChildItem -LiteralPath $FxCRoot -Filter fxc.exe -Recurse -ErrorAction SilentlyContinue |
+    Where-Object { $_.FullName -match '\\x64\\fxc\.exe$' } |
+    Sort-Object -Property FullName -Descending | Select-Object -First 1
+if(-not$FxC){throw 'fxc.exe was not found.'}
+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}
+    & $FxC.FullName /nologo /T cs_5_0 /E $Entry /O3 /Gis /Ges /Fo $OutputPath $ShaderSource
+    $BuildExit=$LASTEXITCODE
+    if($BuildExit-eq-1073741502-and$Attempt-eq0){continue}
+    if($BuildExit-eq-1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
+    if($BuildExit-ne0){throw "fxc failed for $Entry with exit code $BuildExit"}
+    break
+}
+Write-Output "STEP14_SHADER_BUILD entry=$Entry output=$OutputPath exit=0"
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h b/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h
index 3b52a79..0a21685 100644
--- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h
@@ -15,7 +15,16 @@ namespace VoxelForgeGpuFarViewer
 {
     inline constexpr std::uint32_t PooledOutputVertexCapacityPerTile=19456;
     struct CpuVertex { float position[3]{}, normal[3]{}; };
-    struct CpuMesh { std::vector<CpuVertex> vertices; std::vector<std::uint32_t> indices; };
+    struct CpuMesh
+    {
+        std::vector<CpuVertex> vertices;
+        std::vector<std::uint32_t> indices;
+        std::uint32_t surfaceVertexCount=0;
+        std::uint32_t connectorPanelVertexCount=0;
+        std::uint32_t skirtVertexCount=0;
+        bool sequentialIndices=false;
+        bool indirectArgsValid=false;
+    };
     struct GpuTimingQueries
     {
         ID3D11Query* start=nullptr;
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp b/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp
index 3bcb3b3..d459a5c 100644
--- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp
@@ -2352,7 +2352,8 @@ namespace
     }
 
     std::vector<FarSelection> SelectFarTiles(World& world,const WorldConfiguration& config,
-                                             Vec3d camera,bool wholeFar,bool wholeNear)
+                                             Vec3d camera,bool wholeFar,bool wholeNear,
+                                             bool findThinFeature)
     {
         std::vector<FarSelection> result;
         const auto desired=GetFullDesired(world,config,camera);
@@ -2386,6 +2387,43 @@ namespace
                 static_cast<unsigned long long>(perLod[1]));
             return result;
         }
+        if(findThinFeature)
+        {
+            std::uint64_t examined=0,eligible=0;
+            for(const TileKey& key:desired)
+            {
+                if(key.lod<std::max(config.fullResolutionLodCount,2)||key.lod>lastLod)continue;
+                ++examined;
+                const auto origin=TileOrigin(key);
+                const std::int32_t cells=std::max(1,config.coarseTileCells);
+                const std::int32_t step=std::max(1,(32<<key.lod)/cells);
+                if(world.classifyBox({origin[0],origin[1],origin[2]},step,cells)!=TileClass::mixed)
+                    continue;
+                ++eligible;
+                std::string error;
+                const auto exportStart=Clock::now();
+                auto input=world.exportFarFieldTile(key,error);
+                const double exportPreparationMs=std::chrono::duration<double,std::milli>(
+                    Clock::now()-exportStart).count();
+                if(!input)
+                {
+                    if(error.find("near analytic/detail region")!=std::string::npos)continue;
+                    throw std::runtime_error("thin-feature search far field export failed: "+error);
+                }
+                if(input->protectedThinSampleCount==0)continue;
+                std::printf("[ThinFeatureTileSelected] key=(%d,%d,%d) lod=%d protected_thin_samples=%llu examined=%llu eligible=%llu\n",
+                    key.coordinate.x,key.coordinate.y,key.coordinate.z,key.lod,
+                    static_cast<unsigned long long>(input->protectedThinSampleCount),
+                    static_cast<unsigned long long>(examined),
+                    static_cast<unsigned long long>(eligible));
+                result.push_back({key,std::move(*input),exportPreparationMs});
+                return result;
+            }
+            std::printf("[ThinFeatureSearch] examined=%llu eligible=%llu result=no_protected_thin_samples_found\n",
+                static_cast<unsigned long long>(examined),
+                static_cast<unsigned long long>(eligible));
+            return result;
+        }
         for(int32_t lod=std::max(config.fullResolutionLodCount,2);lod<=lastLod;++lod)
         {
             std::vector<TileKey> candidates;
@@ -2978,6 +3016,7 @@ namespace
         const std::filesystem::path shaderPath = root / "Research" / "OwnerShapeWorld" / "gpu_first_port" / "VoxelForgeGpuMesher.hlsl";
         bool wholeFar=false;
         bool wholeNear=false;
+        bool findThinFeature=false;
         bool profileGpuKernels=false;
         bool skirtAuditOnly=false;
         int clipRadiusOverride=-1;
@@ -2990,6 +3029,7 @@ namespace
             if(std::string_view(argv[i])=="--whole-near")wholeNear=true;
             if(std::string_view(argv[i])=="--profile-kernels")profileGpuKernels=true;
             if(std::string_view(argv[i])=="--skirt-audit")skirtAuditOnly=true;
+            if(std::string_view(argv[i])=="--find-thin-feature")findThinFeature=true;
             const std::string_view arg=argv[i];
             if(arg.starts_with("--tile-limit="))
                 tileLimit=std::max(0,std::atoi(std::string(arg.substr(13)).c_str()));
@@ -3068,7 +3108,8 @@ namespace
             noiseWarp.maximum, static_cast<unsigned long long>(noiseWarp.differingBits));
 
         std::vector<FarSelection> selected = skirtAuditOnly
-            ?std::vector<FarSelection>{}:SelectFarTiles(world, config, camera,wholeFar,wholeNear);
+            ?std::vector<FarSelection>{}:SelectFarTiles(world, config, camera,wholeFar,wholeNear,
+                findThinFeature);
         if(skirtAuditOnly)
         {
             if(!config.generateSkirts)
@@ -3411,6 +3452,12 @@ namespace VoxelForgeGpuFarViewer
                 auditReadback!=nullptr,auditReadback!=nullptr,outputSlot,true,captureGpuTiming,batchTimingMode);
             if(auditReadback)
             {
+                auditReadback->surfaceVertexCount=static_cast<std::uint32_t>(result.drawArgs[0]);
+                auditReadback->connectorPanelVertexCount=static_cast<std::uint32_t>(
+                    result.vertices.size()-std::min<std::size_t>(result.vertices.size(),result.drawArgs[0]));
+                auditReadback->skirtVertexCount=static_cast<std::uint32_t>(result.skirtVertices.size());
+                auditReadback->sequentialIndices=result.sequentialIndices;
+                auditReadback->indirectArgsValid=result.indexedIndirectArgsValid;
                 const std::size_t completeGpuVertexCount=result.vertices.size()+result.skirtVertices.size();
                 auditReadback->vertices.resize(completeGpuVertexCount);
                 for(std::size_t i=0;i<result.vertices.size();++i)
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarViewer.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarViewer.ps1
index 9b9fec4..c6b7a10 100644
--- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarViewer.ps1
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarViewer.ps1
@@ -4,7 +4,8 @@ param(
     [string]$LogStem='gpu_far_viewer',
     [int]$ClipRadius=3,
     [string]$CaptureName='',
-    [string]$ExecutableName='VoxelForgeGpuFarHybridViewer.exe'
+    [string]$ExecutableName='VoxelForgeGpuFarHybridViewer.exe',
+    [switch]$WarmupBeforeWorkers
 )
 
 $ErrorActionPreference='Stop'
@@ -23,6 +24,7 @@ $Arguments=@("--clip-radius=$ClipRadius",'--workers=8','--quit-after-clipmap')
 if($Mode-eq'gpu'){$Arguments+='--gpu-far'}
 elseif($Mode-eq'cpu'){$Arguments+='--cpu-far'}
 else{$Arguments+='--gpu-far-audit'}
+if($Mode-eq'gpu'-and$WarmupBeforeWorkers){$Arguments+='--gpu-far-warmup-before-workers'}
 if($CaptureName)
 {
     if([IO.Path]::GetFileName($CaptureName)-ne$CaptureName-or[IO.Path]::GetExtension($CaptureName)-ne'.png')
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15CaptureMatrix.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15CaptureMatrix.ps1
index 9bcd667..1e66c25 100644
--- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15CaptureMatrix.ps1
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15CaptureMatrix.ps1
@@ -1,5 +1,9 @@
 [CmdletBinding()]
-param([string]$ExecutableName='VoxelForgeGpuFieldStep15Timeline.exe',[string]$Prefix='step15')
+param(
+    [string]$ExecutableName='VoxelForgeGpuFieldStep15Timeline.exe',
+    [string]$Prefix='step15',
+    [switch]$GpuWarmupBeforeWorkers
+)
 
 $ErrorActionPreference='Stop'
 $HostProcess=Get-Process -Id $PID
@@ -32,6 +36,7 @@ foreach($View in $Views)
         $Arguments=@('--clip-radius=3','--workers=8','--quit-after-clipmap','--capture-after-clipmap',
             "--camera=$($View.Camera)","--$Mode-far",
             "--capture=Research/OwnerShapeWorld/gpu_field_port/$CaptureName")
+        if($Mode-eq'gpu'-and$GpuWarmupBeforeWorkers){$Arguments+='--gpu-far-warmup-before-workers'}
         $StdoutPath=Join-Path $PSScriptRoot ($Stem+'.stdout.log')
         $StderrPath=Join-Path $PSScriptRoot ($Stem+'.stderr.log')
         $CapturePath=Join-Path $PSScriptRoot $CaptureName
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16CpuDrainControl.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16CpuDrainControl.ps1
new file mode 100644
index 0000000..08008ea
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16CpuDrainControl.ps1
@@ -0,0 +1,68 @@
+[CmdletBinding()]
+param([string]$ExecutableName='VoxelForgeGpuFieldStep16FairRace.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=@(
+    @{Drain=$true;Name='step16_cpu_drain_A'},
+    @{Drain=$false;Name='step16_cpu_paced_A'},
+    @{Drain=$false;Name='step16_cpu_paced_B'},
+    @{Drain=$true;Name='step16_cpu_drain_B'}
+)
+$EightGB=8L*1024L*1024L*1024L
+foreach($Run in $Runs)
+{
+    $Arguments=@('--clip-radius=3','--workers=8','--quit-after-clipmap','--gpu-far-timeline','--cpu-far')
+    if(-not$Run.Drain){$Arguments+='--gpu-far-no-cold-drain'}
+    $StdoutPath=Join-Path $PSScriptRoot ($Run.Name+'.stdout.log')
+    $StderrPath=Join-Path $PSScriptRoot ($Run.Name+'.stderr.log')
+    Write-Output "STEP16_CPU_CONTROL_RUN name=$($Run.Name) drain=$(if($Run.Drain){'unlimited_before_ready'}else{'paced_4_batches_or_2ms'}) 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()
+        $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
+        $Ready=[regex]::Match($LogText,'(?m)^\[InitialClipmapReady\] ([0-9.]+) s')
+        $Presented=[regex]::Match($LogText,'(?m)^\[WholeViewPresented\] at_ms=([0-9.]+) present_ms=([0-9.]+) sync=(\d+) hr=0x([0-9A-Fa-f]{8})')
+        $Frames=[regex]::Matches($LogText,'(?m)^\[ColdFrame\] stage=begin frame=(\d+)')
+        if(-not$Ready.Success-or-not$Presented.Success){throw "Readiness/present marker missing in $($Run.Name); exit=$ViewerExitCode"}
+        $FrameCount=if($Frames.Count){[int]$Frames[$Frames.Count-1].Groups[1].Value}else{0}
+        Write-Output "STEP16_CPU_CONTROL_RESULT name=$($Run.Name) drain=$(if($Run.Drain){'unlimited_before_ready'}else{'paced_4_batches_or_2ms'}) ready_internal_ms=$([math]::Round([double]$Ready.Groups[1].Value*1000,3)) whole_view_present_internal_ms=$($Presented.Groups[1].Value) present_ms=$($Presented.Groups[2].Value) present_hr=0x$($Presented.Groups[4].Value.ToUpperInvariant()) cold_frames=$FrameCount host_process_ms=$([math]::Round($HostRunClock.Elapsed.TotalMilliseconds,3)) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) exit=$ViewerExitCode stdout=$StdoutPath stderr=$StderrPath"
+        if($ViewerExitCode-ne0){throw "Viewer returned $ViewerExitCode`: $($Run.Name)"}
+        break
+    }
+}
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16CpuPresentControl.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16CpuPresentControl.ps1
new file mode 100644
index 0000000..df3d3c9
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16CpuPresentControl.ps1
@@ -0,0 +1,93 @@
+[CmdletBinding()]
+param([string]$ExecutableName='VoxelForgeGpuFieldStep16FairRace.exe')
+
+$ErrorActionPreference='Stop'
+$HostProcess=Get-Process -Id $PID
+$HostProcess.PriorityClass='BelowNormal'
+$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+if(-not('Step16WindowControl' -as [type]))
+{
+    Add-Type @'
+using System;
+using System.Runtime.InteropServices;
+public static class Step16WindowControl
+{
+    [DllImport("user32.dll")] public static extern bool ShowWindow(IntPtr hWnd, int nCmdShow);
+    [DllImport("user32.dll")] public static extern bool SetForegroundWindow(IntPtr hWnd);
+}
+'@
+}
+$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"}
+
+# Same paced CPU flags and workload; only viewer startup visibility changes.
+$Runs=@(
+    @{WindowStyle='Hidden';Name='step16_cpu_paced_hidden_A'},
+    @{WindowStyle='Normal';Name='step16_cpu_paced_visible_A'},
+    @{WindowStyle='Normal';Name='step16_cpu_paced_visible_B'},
+    @{WindowStyle='Hidden';Name='step16_cpu_paced_hidden_B'}
+)
+$EightGB=8L*1024L*1024L*1024L
+foreach($Run in $Runs)
+{
+    $Arguments=@('--clip-radius=3','--workers=8','--quit-after-clipmap','--gpu-far-timeline','--cpu-far','--gpu-far-no-cold-drain')
+    $StdoutPath=Join-Path $PSScriptRoot ($Run.Name+'.stdout.log')
+    $StderrPath=Join-Path $PSScriptRoot ($Run.Name+'.stderr.log')
+    Write-Output "STEP16_PRESENT_CONTROL_RUN name=$($Run.Name) startup_window=$($Run.WindowStyle) 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 $Run.WindowStyle -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
+        $PeakPrivate=0L;$OverMemoryLimit=$false;$ForegroundRaised=$false
+        try
+        {
+            while(-not$Child.HasExited)
+            {
+                $Child.Refresh();if($Child.HasExited){break}
+                $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+                if($Run.WindowStyle-eq'Normal'-and-not$ForegroundRaised)
+                {
+                    $WindowHandle=$Child.MainWindowHandle
+                    if($WindowHandle-ne[IntPtr]::Zero)
+                    {
+                        [void][Step16WindowControl]::ShowWindow($WindowHandle,5)
+                        $ForegroundRaised=[Step16WindowControl]::SetForegroundWindow($WindowHandle)
+                    }
+                }
+                $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()
+        $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
+        $Ready=[regex]::Match($LogText,'(?m)^\[InitialClipmapReady\] ([0-9.]+) s')
+        $Presented=[regex]::Match($LogText,'(?m)^\[WholeViewPresented\] at_ms=([0-9.]+) present_ms=([0-9.]+) sync=(\d+) hr=0x([0-9A-Fa-f]{8})')
+        $Shape=[regex]::Match($LogText,'(?m)^\[ColdTimeline\] stage=shape_compile end_ms=[0-9.]+ duration_ms=([0-9.]+)')
+        $Frames=[regex]::Matches($LogText,'(?m)^\[ColdFrame\] stage=begin frame=(\d+)')
+        if(-not$Ready.Success-or-not$Presented.Success-or-not$Shape.Success)
+        {throw "Required timing/present marker missing in $($Run.Name); exit=$ViewerExitCode"}
+        $PresentHR=$Presented.Groups[4].Value.ToUpperInvariant()
+        $PresentState=if($PresentHR-eq'087A0001'){'occluded'}elseif($PresentHR-eq'00000000'){'visible_success'}else{'other'}
+        $FrameCount=if($Frames.Count){[int]$Frames[$Frames.Count-1].Groups[1].Value}else{0}
+        Write-Output "STEP16_PRESENT_CONTROL_RESULT name=$($Run.Name) startup_window=$($Run.WindowStyle) foreground_raised=$ForegroundRaised ready_internal_ms=$([math]::Round([double]$Ready.Groups[1].Value*1000,3)) whole_view_present_internal_ms=$($Presented.Groups[1].Value) present_ms=$($Presented.Groups[2].Value) present_sync=$($Presented.Groups[3].Value) present_hr=0x$PresentHR present_state=$PresentState cold_frames=$FrameCount shape_compile_ms=$($Shape.Groups[1].Value) host_process_ms=$([math]::Round($HostRunClock.Elapsed.TotalMilliseconds,3)) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) exit=$ViewerExitCode stdout=$StdoutPath stderr=$StderrPath"
+        if($ViewerExitCode-ne0){throw "Viewer returned $ViewerExitCode`: $($Run.Name)"}
+        break
+    }
+}
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16FairRace.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16FairRace.ps1
new file mode 100644
index 0000000..fd79f43
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16FairRace.ps1
@@ -0,0 +1,84 @@
+[CmdletBinding()]
+param([string]$ExecutableName='VoxelForgeGpuFieldStep16FairRace.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"}
+
+# Counterbalance which mode goes first in each pair while keeping launches interleaved.
+$Runs=@(
+    @{Mode='gpu';Name='step16_race_gpu_A'},
+    @{Mode='cpu';Name='step16_race_cpu_A'},
+    @{Mode='cpu';Name='step16_race_cpu_B'},
+    @{Mode='gpu';Name='step16_race_gpu_B'},
+    @{Mode='gpu';Name='step16_race_gpu_C'},
+    @{Mode='cpu';Name='step16_race_cpu_C'},
+    @{Mode='cpu';Name='step16_race_cpu_D'},
+    @{Mode='gpu';Name='step16_race_gpu_D'}
+)
+$EightGB=8L*1024L*1024L*1024L
+$ModeCounts=@{cpu=0;gpu=0}
+foreach($Run in $Runs)
+{
+    $ModeCounts[$Run.Mode]++
+    $LaunchClass=if($ModeCounts[$Run.Mode]-eq1){'first_mode_launch'}else{'subsequent_mode_launch'}
+    $Arguments=@('--clip-radius=3','--workers=8','--quit-after-clipmap','--gpu-far-timeline',"--$($Run.Mode)-far")
+    $StdoutPath=Join-Path $PSScriptRoot ($Run.Name+'.stdout.log')
+    $StderrPath=Join-Path $PSScriptRoot ($Run.Name+'.stderr.log')
+    Write-Output "STEP16_RUN mode=$($Run.Mode) name=$($Run.Name) launch_class=$LaunchClass 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()
+        $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
+        $Ready=[regex]::Match($LogText,'(?m)^\[InitialClipmapReady\] ([0-9.]+) s')
+        $Presented=[regex]::Match($LogText,'(?m)^\[WholeViewPresented\] at_ms=([0-9.]+) present_ms=([0-9.]+) sync=(\d+) hr=0x([0-9A-Fa-f]{8})')
+        $Shape=[regex]::Match($LogText,'(?m)^\[ColdTimeline\] stage=shape_compile end_ms=[0-9.]+ duration_ms=([0-9.]+)')
+        $Frames=[regex]::Matches($LogText,'(?m)^\[ColdFrame\] stage=begin frame=(\d+)')
+        $Memory=[regex]::Match($LogText,'(?m)^\[GpuFarMemory\] dxgi_local_sampled_peak_mb=([0-9.]+)')
+        if(-not$Ready.Success-or-not$Presented.Success-or-not$Shape.Success)
+        {throw "Required timing/present marker missing in $($Run.Name); exit=$ViewerExitCode"}
+        $PresentHR=$Presented.Groups[4].Value.ToUpperInvariant()
+        $PresentState=if($PresentHR-eq'087A0001'){'occluded'}elseif($PresentHR-eq'00000000'){'visible_success'}else{'other'}
+        $DxgiMB=if($Memory.Success){[double]$Memory.Groups[1].Value}else{$null}
+        $CombinedGB=$PeakPrivate/1GB
+        if($null-ne$DxgiMB){$CombinedGB=($PeakPrivate+$DxgiMB*1MB)/1GB}
+        if($CombinedGB-ge8.0){throw "Viewer exceeded conservative private + DXGI-local 8 GiB ceiling: $($Run.Name) ($CombinedGB GiB)"}
+        $FrameCount=if($Frames.Count){[int]$Frames[$Frames.Count-1].Groups[1].Value}else{0}
+        Write-Output "STEP16_RESULT mode=$($Run.Mode) name=$($Run.Name) launch_class=$LaunchClass ready_internal_ms=$([math]::Round([double]$Ready.Groups[1].Value*1000,3)) whole_view_present_internal_ms=$($Presented.Groups[1].Value) present_ms=$($Presented.Groups[2].Value) present_sync=$($Presented.Groups[3].Value) present_hr=0x$PresentHR present_state=$PresentState cold_frames=$FrameCount shape_compile_ms=$($Shape.Groups[1].Value) host_process_ms=$([math]::Round($HostRunClock.Elapsed.TotalMilliseconds,3)) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) dxgi_local_peak_mb=$(if($null-ne$DxgiMB){$DxgiMB}else{'not_reported'}) conservative_private_plus_dxgi_gb=$([math]::Round($CombinedGB,2)) exit=$ViewerExitCode stdout=$StdoutPath stderr=$StderrPath"
+        if($ViewerExitCode-ne0){throw "Viewer returned $ViewerExitCode`: $($Run.Name)"}
+        break
+    }
+}
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16GpuWarmupControl.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16GpuWarmupControl.ps1
new file mode 100644
index 0000000..91bad33
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep16GpuWarmupControl.ps1
@@ -0,0 +1,75 @@
+[CmdletBinding()]
+param([string]$ExecutableName='VoxelForgeGpuFieldStep16FairRace.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"}
+
+# Counterbalance warm-up and no-warm-up launches; all other arguments stay fixed.
+$Runs=@(
+    @{Prewarm=$true;Name='step16_gpu_warmup_A'},
+    @{Prewarm=$false;Name='step16_gpu_no_prewarm_A'},
+    @{Prewarm=$false;Name='step16_gpu_no_prewarm_B'},
+    @{Prewarm=$true;Name='step16_gpu_warmup_B'}
+)
+$EightGB=8L*1024L*1024L*1024L
+foreach($Run in $Runs)
+{
+    $Arguments=@('--clip-radius=3','--workers=8','--quit-after-clipmap','--gpu-far-timeline','--gpu-far')
+    if(-not$Run.Prewarm){$Arguments+='--gpu-far-no-prewarm'}
+    $StdoutPath=Join-Path $PSScriptRoot ($Run.Name+'.stdout.log')
+    $StderrPath=Join-Path $PSScriptRoot ($Run.Name+'.stderr.log')
+    Write-Output "STEP16_GPU_WARMUP_CONTROL_RUN name=$($Run.Name) prewarm=$($Run.Prewarm) 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()
+        $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
+        $Ready=[regex]::Match($LogText,'(?m)^\[InitialClipmapReady\] ([0-9.]+) s')
+        $Presented=[regex]::Match($LogText,'(?m)^\[WholeViewPresented\] at_ms=([0-9.]+) present_ms=([0-9.]+) sync=(\d+) hr=0x([0-9A-Fa-f]{8})')
+        $Shape=[regex]::Match($LogText,'(?m)^\[ColdTimeline\] stage=shape_compile end_ms=[0-9.]+ duration_ms=([0-9.]+)')
+        $WarmupQueue=[regex]::Match($LogText,'(?m)^\[GpuShaderWarmup\] dispatched=.* queue_ms=([0-9.]+)')
+        $FirstPresent=[regex]::Match($LogText,'(?m)^\[PresentProfile\] ordinal=1 .* call_ms=([0-9.]+)')
+        $Memory=[regex]::Match($LogText,'(?m)^\[GpuFarMemory\] dxgi_local_sampled_peak_mb=([0-9.]+)')
+        if(-not$Ready.Success-or-not$Presented.Success-or-not$Shape.Success)
+        {throw "Required timing/present marker missing in $($Run.Name); exit=$ViewerExitCode"}
+        $DxgiMB=if($Memory.Success){[double]$Memory.Groups[1].Value}else{0.0}
+        $CombinedGB=($PeakPrivate+$DxgiMB*1MB)/1GB
+        if($CombinedGB-ge8.0){throw "Viewer exceeded conservative private + DXGI-local 8 GiB ceiling: $($Run.Name) ($CombinedGB GiB)"}
+        Write-Output "STEP16_GPU_WARMUP_CONTROL_RESULT name=$($Run.Name) prewarm=$($Run.Prewarm) ready_internal_ms=$([math]::Round([double]$Ready.Groups[1].Value*1000,3)) whole_view_present_internal_ms=$($Presented.Groups[1].Value) present_ms=$($Presented.Groups[2].Value) present_hr=0x$($Presented.Groups[4].Value.ToUpperInvariant()) shape_compile_ms=$($Shape.Groups[1].Value) warmup_queue_ms=$(if($WarmupQueue.Success){$WarmupQueue.Groups[1].Value}else{'not_reported'}) first_present_call_ms=$(if($FirstPresent.Success){$FirstPresent.Groups[1].Value}else{'not_reported'}) host_process_ms=$([math]::Round($HostRunClock.Elapsed.TotalMilliseconds,3)) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) dxgi_local_peak_mb=$DxgiMB conservative_private_plus_dxgi_gb=$([math]::Round($CombinedGB,2)) exit=$ViewerExitCode stdout=$StdoutPath stderr=$StderrPath"
+        if($ViewerExitCode-ne0){throw "Viewer returned $ViewerExitCode`: $($Run.Name)"}
+        break
+    }
+}
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep17TileControls.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep17TileControls.ps1
new file mode 100644
index 0000000..f773ad8
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep17TileControls.ps1
@@ -0,0 +1,117 @@
+[CmdletBinding()]
+param(
+    [string]$ExecutableName='VoxelForgeGpuFieldStep17.exe',
+    [string]$Prefix='step17_tilecontrol',
+    [ValidateSet('all','cpu','gpu','gpu_direct_draw','gpu_readback_copy','gpu_meter_copy','gpu_cpu_mesh_copy','gpu_depth_disabled','gpu_skirts_off','gpu_connectors_off','gpu_skirts_connectors_off')]
+    [string]$OnlyControl='all'
+)
+
+$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"}
+
+$Controls=@(
+    @{Name='cpu';Flags=@('--cpu-far')},
+    @{Name='gpu';Flags=@('--gpu-far-warmup-before-workers')},
+    @{Name='gpu_direct_draw';Flags=@('--gpu-far-warmup-before-workers','--gpu-far-direct-draw-audit')},
+    @{Name='gpu_readback_copy';Flags=@('--gpu-far-warmup-before-workers','--gpu-far-readback-copy-draw-audit')},
+    @{Name='gpu_meter_copy';Flags=@('--gpu-far-warmup-before-workers','--gpu-far-meter-copy-draw-audit')},
+    @{Name='gpu_cpu_mesh_copy';Flags=@('--gpu-far-warmup-before-workers','--gpu-far-cpu-mesh-draw-audit')},
+    @{Name='gpu_depth_disabled';Flags=@('--gpu-far-warmup-before-workers','--gpu-far-depth-disabled-audit')},
+    @{Name='gpu_skirts_off';Flags=@('--gpu-far-warmup-before-workers','--gpu-skirts-disabled')},
+    @{Name='gpu_connectors_off';Flags=@('--gpu-far-warmup-before-workers','--gpu-connectors-disabled')},
+    @{Name='gpu_skirts_connectors_off';Flags=@('--gpu-far-warmup-before-workers','--gpu-skirts-disabled','--gpu-connectors-disabled')}
+)
+if($OnlyControl-ne'all'){$Controls=@($Controls|Where-Object{$_.Name-eq$OnlyControl})}
+$EightGB=8L*1024L*1024L*1024L
+foreach($Control in $Controls)
+{
+    $Stem="${Prefix}_$($Control.Name)"
+    $CaptureName="$Stem.png"
+    $CapturePath=Join-Path $PSScriptRoot $CaptureName
+    $StdoutPath=Join-Path $PSScriptRoot ($Stem+'.stdout.log')
+    $StderrPath=Join-Path $PSScriptRoot ($Stem+'.stderr.log')
+    $Arguments=@('--clip-radius=3','--workers=8','--quit-after-clipmap','--capture-after-clipmap',
+        '--camera=500,0,0,0,0',
+        '--draw-only-tile=6,2,0,3','--gpu-skirt-draw-audit',"--capture=Research/OwnerShapeWorld/gpu_field_port/$CaptureName")
+    $Arguments+=$Control.Flags
+    Write-Output "STEP17_TILE_CONTROL name=$($Control.Name) priority=BelowNormal 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 "Tile control stopped at the 8 GiB private-memory ceiling: $($Control.Name)"}
+        $LogText=Get-Content -LiteralPath $StdoutPath -Raw
+        if(-not$LogText.Contains('[InitialClipmapReady]')-or-not(Test-Path -LiteralPath $CapturePath -PathType Leaf))
+        {throw "Full-residency tile control capture missing: $($Control.Name); exit=$ChildExitCode"}
+        if($Control.Name-ne'cpu')
+        {
+            if(-not$LogText.Contains('key=(6,2,0,lod3) route=gpu_eligible'))
+            {throw "Target tile was not directly GPU eligible: $($Control.Name)"}
+            if(-not$LogText.Contains('[GpuSkirtDrawArgs] key=(6,2,0,lod3)'))
+            {throw "Target draw arguments were not audited: $($Control.Name)"}
+            $PooledOutputAudit=[regex]::Match($LogText,'(?m)^\[GpuPooledOutputReadback\] key=\(6,2,0,lod3\).*sequential_indices=(\d+) indirect_args_valid=(\d+).*status=(PASS|FAIL)')
+            if(-not$PooledOutputAudit.Success-or$PooledOutputAudit.Groups[1].Value-ne'1'-or$PooledOutputAudit.Groups[2].Value-ne'1'-or$PooledOutputAudit.Groups[3].Value-ne'PASS')
+            {throw "Target pooled GPU indices or indirect arguments failed readback audit: $($Control.Name)"}
+        }
+        $Memory=[regex]::Match($LogText,'(?m)^\[GpuFarMemory\] dxgi_local_sampled_peak_mb=([0-9.]+)')
+        $DxgiMB=if($Memory.Success){[double]$Memory.Groups[1].Value}else{0.0}
+        $CombinedGB=($PeakPrivate+$DxgiMB*1MB)/1GB
+        if($CombinedGB-ge8.0){throw "Tile control exceeded private + DXGI-local 8 GiB: $($Control.Name) ($CombinedGB GiB)"}
+        $ArgsAudit=[regex]::Match($LogText,'(?m)^\[GpuSkirtDrawArgs\] key=\(6,2,0,lod3\).*surface_indices=(\d+).*surface_start_index=(\d+).*expected_surface_start_index=(\d+).*skirt_panel_indices=(\d+).*connector_start_index=(\d+).*expected_connector_start_index=(\d+)')
+        if($Control.Name-ne'cpu')
+        {
+            $ArgsSummary=[regex]::Match($LogText,'(?m)^\[GpuSkirtDrawArgsSummary\].*status=(PASS|FAIL)')
+            if(-not$ArgsAudit.Success-or-not$ArgsSummary.Success-or$ArgsSummary.Groups[1].Value-ne'PASS')
+            {throw "Target indirect arguments failed audit: $($Control.Name)"}
+            if($ArgsAudit.Groups[2].Value-ne$ArgsAudit.Groups[3].Value-or$ArgsAudit.Groups[5].Value-ne$ArgsAudit.Groups[6].Value)
+            {throw "Target pooled draw starts mismatch: $($Control.Name)"}
+            $TriangleAudit=[regex]::Match($LogText,'(?m)^\[GpuCpuTriangleAudit\] key=\(6,2,0,lod3\).*cpu_surface_triangles=(\d+) gpu_surface_triangles=(\d+).*unmatched_cpu=(\d+) unmatched_gpu=(\d+)')
+            if(-not$TriangleAudit.Success-or$TriangleAudit.Groups[1].Value-ne$TriangleAudit.Groups[2].Value-or$TriangleAudit.Groups[3].Value-ne'0'-or$TriangleAudit.Groups[4].Value-ne'0')
+            {throw "Target CPU/GPU surface triangle audit failed: $($Control.Name)"}
+            if($Control.Name-eq'gpu_direct_draw'-and-not$LogText.Contains('[GpuDirectDrawAudit] key=(6,2,0,lod3)'))
+            {throw 'Direct GPU draw control did not use the audited direct path.'}
+            if(($Control.Name -eq 'gpu_readback_copy' -or $Control.Name -eq 'gpu_meter_copy') -and
+                -not $LogText.Contains('[GpuReadbackCopyDrawAudit] key=(6,2,0,lod3)'))
+            {throw 'Readback-copy GPU draw control did not create an unpooled draw buffer.'}
+            if($Control.Name -eq 'gpu_cpu_mesh_copy' -and
+                -not $LogText.Contains('[GpuCpuMeshDrawAudit] key=(6,2,0,lod3)'))
+            {throw 'CPU-mesh GPU draw control did not create its independent draw buffer.'}
+            if($Control.Name -eq 'gpu_depth_disabled' -and
+                -not $LogText.Contains('[GpuFarDepthDisabledAudit] key=(6,2,0,lod3) depth_test=disabled'))
+            {throw 'Depth-disabled GPU draw control did not disable depth testing on the target tile.'}
+        }
+        Write-Output "STEP17_TILE_CONTROL_DONE name=$($Control.Name) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) dxgi_local_peak_mb=$DxgiMB conservative_private_plus_dxgi_gb=$([math]::Round($CombinedGB,2)) target_draw_args=$(if($ArgsAudit.Success){$ArgsAudit.Value}else{'not_applicable'}) capture=$CapturePath exit=$ChildExitCode stdout=$StdoutPath stderr=$StderrPath"
+        if($ChildExitCode-ne0){throw "Viewer returned $ChildExitCode`: $($Control.Name)"}
+        break
+    }
+}
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/STEP16_REPORT.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/STEP16_REPORT.md
new file mode 100644
index 0000000..1e2c406
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/STEP16_REPORT.md
@@ -0,0 +1,70 @@
+# GPU field port — Step 16 report
+
+## Result
+
+The CPU viewer now drains all dirty CPU upload batches before the first fully resident view. Once `[InitialClipmapReady]` is reported, the existing four-batch / 2 ms per-frame upload budget applies again. In the interleaved Step 16 race, the hybrid still loses: median internal readiness was **3.152 s GPU vs 1.354 s CPU**, a **1.799 s CPU lead**. Median host process time was **3.341 s GPU vs 1.484 s CPU**.
+
+The upload cap is no longer the CPU cold-start gate. In this machine's occluded-window runs, worker completion was the final gate, so unlimited draining did not materially improve the total CPU time over the paced control.
+
+## CPU cold drain
+
+`UploadDirtyBatches` now drains without the per-frame caps until the initial clipmap is ready, in both CPU and hybrid modes. `--gpu-far-no-cold-drain` remains the paced control. The normal four-batch / 2 ms limit resumes after readiness for VR frame pacing. See [StandaloneViewer.cpp](../../../Source/StandaloneViewer.cpp) and [CPU drain control runner](RunStep16CpuDrainControl.ps1).
+
+The control logs show the cap being bypassed and where the remaining time goes:
+
+| CPU run | Upload mode | Max uploaded batches in a frame | Last CPU job finish | Last upload | Ready |
+|---|---|---:|---:|---:|---:|
+| Drain A | Unlimited | 45 | 1,428.551 ms, frame 503 | 1,433.374 ms, frame 504 | 1,433 ms |
+| Drain B | Unlimited | 63 | 1,364.572 ms, frame 315 | 1,368.638 ms, frame 316 | 1,369 ms |
+| Paced A | 4 batches / 2 ms | 4 | 1,381.804 ms, frame 154 | 1,388.175 ms, frame 156 | 1,388 ms |
+| Paced B | 4 batches / 2 ms | 4 | 1,381.767 ms, frame 151 | 1,389.015 ms, frame 154 | 1,389 ms |
+
+The cold-drain path uploads all currently dirty batches, with up to 45–63 batch rebuilds in one frame. The final worker result then reaches the renderer in the next draw pass, about 4–9 ms later. The paced control happened to finish producing dirty work gradually on this machine, so its total readiness was similar; it never had more than four uploaded batches in a frame. Cold readiness still waits for CPU generation and the next renderer pass, but it no longer waits for a multi-frame upload-budget tail.
+
+## Fair race
+
+The race used one binary, `VoxelForgeGpuFieldStep16FairRace.exe`, and identical shared flags for all eight runs: `--clip-radius=3 --workers=8 --quit-after-clipmap --gpu-far-timeline`. Only `--cpu-far` versus `--gpu-far` differed. Each process ran at BelowNormal priority with affinity `0xFF00`; launches were interleaved and counterbalanced. Each run reports the viewer's internal ready time and the host stopwatch time through process exit. “First process” means the first viewer launch for that mode in this race; no OS or driver cache was cleared.
+
+| Run | Launch class | Ready (internal) | Host exit | Cold frames | Shape compile | Whole-view present |
+|---|---|---:|---:|---:|---:|---|
+| GPU A | First GPU process | 3,176 ms | 3,359 ms | 3 | 849 ms | Occluded, 0.372 ms |
+| CPU A | First CPU process | 1,408 ms | 1,558 ms | 205 | 627 ms | Occluded, 0.095 ms |
+| CPU B | Subsequent | 1,357 ms | 1,470 ms | 26 | 597 ms | Occluded, 0.071 ms |
+| GPU B | Subsequent | 3,098 ms | 3,253 ms | 3 | 864 ms | Occluded, 0.377 ms |
+| GPU C | Subsequent | 3,255 ms | 3,439 ms | 3 | 846 ms | Occluded, 0.424 ms |
+| CPU C | Subsequent | 1,349 ms | 1,464 ms | 111 | 584 ms | Occluded, 0.082 ms |
+| CPU D | Subsequent | 1,350 ms | 1,497 ms | 52 | 591 ms | Occluded, 0.101 ms |
+| GPU D | Subsequent | 3,128 ms | 3,322 ms | 4 | 851 ms | Occluded, 0.267 ms |
+
+First process launch for each mode: GPU **3.176 s** vs CPU **1.408 s** internal readiness. Subsequent-launch medians: GPU **3.128 s** vs CPU **1.350 s**. Across all four launches per mode, the medians are **3.152 s GPU vs 1.354 s CPU** internal and **3.341 s GPU vs 1.484 s CPU** host exit. The hybrid does **not** beat CPU in this race.
+
+All eight whole-view presents returned `DXGI_STATUS_OCCLUDED` (`0x087A0001`). The CPU cold-frame ordinal ranged from 26 to 205 while readiness stayed within 59 ms; GPU ordinals were 3–4. The frame counts are in the timeline logs, but are not a direct proxy for elapsed time: CPU frame work varied with arriving worker results, while the GPU path spent most of its interval inside one long pump. The older Step 15 race and `claude_s15` runs did not enable frame telemetry, so their exact frame counts cannot be recovered.
+
+GPU local memory in the race peaked at 1,764.2 MB; combined with peak private memory, the largest sampled total was **4.96 GB**. CPU private memory peaked at **0.88 GB**. CPU mode does not emit the GPU-local sample; its final uploaded geometry payload is about 55 MiB (1,471,806 24-byte vertices plus 5,584,440 32-bit indices).
+
+## Why timings vary
+
+The old CPU shape-compile measurements are close across the two sets of Step 15 logs: **585–586 ms** in the Step 15 race and **589–599 ms** in the `claude_s15` runs. That does not explain the readiness gap between the user's **2.06–2.25 s** CPU runs and the Step 15 race's **1.36–1.42 s** runs.
+
+The present logs show a real difference between those sets. The `claude_s15` CPU and GPU runs returned `S_OK`; CPU startup present calls ranged from 14 ms to 194 ms. The Step 15 race and all Step 16 race runs returned `DXGI_STATUS_OCCLUDED`; whole-view present calls were below 0.4 ms. This presentation difference is consistent with frame-paced work progressing differently, but it was not isolated as the cause: a Step 16 hidden-versus-normal launch control still returned occluded in both cases, so it did not change the renderer's actual presentation state.
+
+For the user's Step 15 GPU outlier, `claude_s15_gpu_a` reported **2.902 s** readiness, **895 ms** shape compilation, **1,226 ms** shader-warmup queue time, and a **322 ms** first present call. The two later GPU runs reported **1.252 / 1.209 s**, **594 / 583 ms** shape compilation, **3.2 / 3.4 ms** warmup queue time, and **1.8 ms** first present calls. Those markers move together, but that run set does not isolate one as the cause.
+
+In the final workspace state, Step 16 shape compilation was **845–864 ms GPU** (median 850 ms) and **584–627 ms CPU** (median 594 ms). I ran a controlled warmup toggle with two launches per setting. Async warmup produced **876–912 ms** shape-compile times; `--gpu-far-no-prewarm` produced **839–865 ms**. Disabling warmup reduced the median compile time by about **42 ms**, but total readiness stayed around **3.1–3.3 s**: warmup median **3.205 s**, no-prewarm median **3.221 s**. The no-prewarm first-present calls were 22 ms, compared with 222–343 ms with warmup, but whole-view readiness did not improve; these runs do not show where the saved present-call time is paid.
+
+The final race's GPU pump wall time was **1.71–1.90 s**; first submission came at **2.82–2.94 s** and completion at **3.09–3.24 s**. The GPU dispatch timestamp ranged from **199 to 400 ms**, with one run's slowest tile at 201 ms. The GPU first-present call ranged from **123 to 300 ms**, although those calls returned the occluded status. These measured stages account for most of the GPU time and vary substantially; the logs do not isolate why the pump or GPU dispatch slowed on some launches.
+
+CPU last-job completion varied from **1,343 to 1,402 ms**, followed by final upload within **6–7 ms**. The shape-compile gap is real, and the no-prewarm control indicates async warmup accounts for a modest part of it. Warmup alone does not explain the roughly 1.8 s race gap: the remaining cost is inside the GPU pump, submission, and dispatch path. The GPU slowdown seen here was present on both the first and subsequent process launches. The Step 15 `2.9 s` first-GPU outlier is therefore not a one-time outlier in this final-state run set; its earlier shorter repeats were not reproduced.
+
+## Correctness and resource checks
+
+- CPU reference W1 and W8 hashes match: density `0xdb2991c73215af4c`, mesh `0x8236c5edd7fc8e4f`, payload `0x4c48fa75a543d58e`, topology `0xbea22e76b34a4dfc`. Walkability passed **5,642/5,642**. Logs: [W1](step16_cpu_reference_w1.stdout.log), [W8](step16_cpu_reference_w8.stdout.log).
+- Whole-near agreement passed on **960 tiles**: zero GPU/WARP topology differences; **229,376** protected-exception cells with zero GPU/WARP sign or topology differences. Whole-far passed on **1,114 tiles** (152 transition tiles skipped): zero GPU/WARP sign differences across tested far-field samples; **127,401** protected-exception cells and **353** thin-feature cells with zero GPU/WARP sign or topology differences. Logs: [whole-near](step16_whole_near_full_agreement.stdout.log), [whole-far](step16_whole_far_agreement.stdout.log).
+- The GPU survival audit passed at **3.11 GB private + 2,043.8 MB DXGI local = 5.10 GB**. Log: [survival audit](step16_gpu_survival.stdout.log).
+- All ten full-residency CPU/GPU capture pairs completed. The 20 PNGs were only read as numeric RGB arrays; no full-size capture was opened in an image viewer. Results are in [capture comparison](step16_capture_comparison.md); the largest sampled GPU capture total was 5.22 GB.
+- Build used `/MP8`; runner processes used BelowNormal and `0xFF00`. Runners allowed one retry on `0xC0000142` and closed every viewer they launched. No commit, push, or stash was made.
+- The build helper regenerated tracked `precompiled/EmitNearCS.cso`; I restored it from `HEAD` before the final race and GPU checks, and it now matches `HEAD`. No system or driver cache was cleared or modified.
+
+## Remaining coverage limitation
+
+Direct-GPU LOD3/4 coverage remains unchanged. The hybrid currently routes all LOD3 and LOD4 tiles to CPU for `remote_capture_pixel_coverage`; this run had 208 LOD3 and 144 LOD4 remote-band fallback tiles and 1,144 CPU fallback tiles total. The `(6,2,0,lod3)` direct-GPU coverage issue remains open; no direct-GPU coverage change was attempted in this step.
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/STEP17_REPORT.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/STEP17_REPORT.md
new file mode 100644
index 0000000..1bd8edc
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/STEP17_REPORT.md
@@ -0,0 +1,42 @@
+# GPU port, step 17
+
+## Status
+
+Goal 2 is fixed and validated on the final r29 build. Goal 3 is measured below. Goal 1 remains open: cold adaptive startup still loses to CPU-only startup in this sandbox, and the cached race was not measured here.
+
+## Goal 1: adaptive startup
+
+Final r29 cold run, `--gpu-far` adaptive mode:
+
+- CPU reference: 1.352 s to the initial clipmap; CPU mesh worker sum 5.046 s.
+- Adaptive GPU mode: 1.867 s to the initial clipmap; CPU mesh worker sum 6.690 s; GPU busy timestamp sum 0 ms.
+- The cold shader queue took 1,087 ms. All 1,434 GPU-eligible tiles were already started on CPU at handoff; the handoff audit passed, but no tile was available for GPU takeover.
+
+Two cache-probe experiments were rejected and removed from the viewer source. The single-shader probe took 404 ms on r27 and 912 ms on r28, slowing startup further. Their logs remain as evidence. No driver or shader cache was read or changed. The owner should run the cached race outside the sandbox.
+
+## Goal 2: LOD3/4 coverage
+
+The broad missing patch came from GPU skirt depth using `Key.StepVoxels()` instead of the exported coarse mesher step. At `(6,2,0,lod3)`, that made the GPU skirt 800 cm deep while the CPU skirt was 1,600 cm deep. The viewer now uses `FarFieldTileInput.step` for skirt depth in [StandaloneViewer.cpp](../../../Source/StandaloneViewer.cpp).
+
+On r29, the isolated `(6,2,0,lod3)` tile passed the GPU readback, indirect-argument, and one-to-one triangle audits: 680 CPU/GPU surface triangles and 64 visible skirt triangles matched; maximum surface vertex error was 0.001892 cm. The CPU and GPU captures were pixel-identical at 1280×800, with zero missing pixels. The numerical report is [step17_final_r29_tilecontrol_comparison.md](step17_final_r29_tilecontrol_comparison.md).
+
+All ten r29 CPU/GPU view pairs completed after full clipmap residency. Eight views had zero CPU-rock-to-GPU-background pixels. `remote_yaw180` and `remote_yaw270` each had two such pixels; neither showed the broad patch. Peak private plus DXGI-local memory across the matrix was 6.17 GB. See [step17_final_r29_capture_comparison.md](step17_final_r29_capture_comparison.md).
+
+The final r29 GPU survival audit passed (`mesh_source=GPU`, required field readback); peak private plus DXGI-local memory was 6.15 GB. The thin-feature selector in [GpuFieldPort.cpp](GpuFieldPort.cpp) found tile `(-1,0,-1,lod3)`: 18 thin-feature cells and 35 protected-exception cells were audited, with zero topology differences against CPU on both hardware GPU and WARP. See `step17_thin_feature_audit.stdout.log`.
+
+CPU reference hashes matched at 1 and 8 workers, with walkability 5,642/5,642:
+
+`0xdb2991c73215af4c, 0x8236c5edd7fc8e4f, 0x4c48fa75a543d58e, 0xbea22e76b34a4dfc`
+
+## Goal 3: timing
+
+| Mode | CPU mesh worker sum | CPU input/export worker sum | GPU busy timestamp sum | Initial clipmap wall |
+|---|---:|---:|---:|---:|
+| CPU-only, r29 baseline | 5.046 s | — | 0 ms | 1.352 s |
+| Adaptive GPU, cold r29 | 6.690 s | 0 s | 0 ms | 1.867 s |
+| GPU hybrid, ten cold full-residency views | 0.830–1.408 s | 0.569–0.971 s | 219.181–226.235 ms | 1.813–2.207 s |
+| GPU survival audit, r29 | 1.166 s | 1.207 s | 793.829 ms | 44.802 s |
+
+The capture-matrix hybrid runs used `--gpu-far-warmup-before-workers` so the captures exercise the GPU path even with a cold sandbox driver. Their GPU timing is a sum of per-dispatch GPU timestamps, not wall time. The survival run includes additional audit/readback work.
+
+All work and artifacts remain under `Standalone/`. No images were opened in a viewer; capture comparisons were numerical. Nothing was committed, pushed, or stashed.
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/EmitNearCS.cso b/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/EmitNearCS.cso
index bcb1211..cc34ff4 100644
Binary files a/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/EmitNearCS.cso and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/EmitNearCS.cso differ
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step16_capture_comparison.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step16_capture_comparison.md
new file mode 100644
index 0000000..ef394a8
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step16_capture_comparison.md
@@ -0,0 +1,16 @@
+# Step16 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.2002 | 2.6637 | 4.2874% | 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.2360 | 3.1110 | 5.5744% | 0.4098% | 107 |
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_final_capture_comparison.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_final_capture_comparison.md
new file mode 100644
index 0000000..f6ba3ad
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_final_capture_comparison.md
@@ -0,0 +1,18 @@
+# Step17_Final 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.
+
+`CPU rock to GPU background` isolates the prior missing-coverage signature. Captures were loaded as RGB arrays only; no full-size image was opened in a viewer.
+
+| View | Size | RGB MAE | RGB RMSE | Changed pixels | Pixels max channel >32 | Max channel | CPU rock to GPU background | GPU rock to CPU background | Missing-pixel bounds |
+|---|---:|---:|---:|---:|---:|---:|---:|---:|---|
+| chamber | 1280x800 | 0.6075 | 1.8638 | 42.6368% | 0.0183% | 123 | 0 | 0 | none |
+| tunnel | 1280x800 | 0.6383 | 3.3799 | 31.7335% | 0.2772% | 130 | 0 | 0 | none |
+| passage | 1280x800 | 1.6092 | 3.6878 | 50.4982% | 0.1733% | 52 | 0 | 0 | none |
+| far | 1280x800 | 0.7910 | 2.1578 | 44.8325% | 0.0198% | 128 | 0 | 0 | none |
+| passage_floor_closeup | 1280x800 | 0.9314 | 2.4800 | 43.0219% | 0.0500% | 100 | 0 | 0 | none |
+| tunnel_bend | 1280x800 | 0.8129 | 2.1599 | 40.9629% | 0.0708% | 55 | 0 | 0 | none |
+| remote_yaw0 | 1280x800 | 0.0416 | 1.0679 | 2.7991% | 0.0305% | 90 | 0 | 0 | none |
+| remote_yaw90 | 1280x800 | 0.0100 | 0.4323 | 0.9746% | 0.0035% | 94 | 0 | 0 | none |
+| remote_yaw180 | 1280x800 | 0.0470 | 1.3091 | 2.9463% | 0.0311% | 102 | 2 | 0 | (415,172)-(834,173) |
+| remote_yaw270 | 1280x800 | 0.0826 | 1.6874 | 5.1076% | 0.0580% | 107 | 2 | 0 | (835,204)-(886,205) |
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_final_r29_capture_comparison.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_final_r29_capture_comparison.md
new file mode 100644
index 0000000..fd9ba4e
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_final_r29_capture_comparison.md
@@ -0,0 +1,18 @@
+# Step17_Final_R29 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.
+
+`CPU rock to GPU background` isolates the prior missing-coverage signature. Captures were loaded as RGB arrays only; no full-size image was opened in a viewer.
+
+| View | Size | RGB MAE | RGB RMSE | Changed pixels | Pixels max channel >32 | Max channel | CPU rock to GPU background | GPU rock to CPU background | Missing-pixel bounds |
+|---|---:|---:|---:|---:|---:|---:|---:|---:|---|
+| chamber | 1280x800 | 0.6075 | 1.8638 | 42.6368% | 0.0183% | 123 | 0 | 0 | none |
+| tunnel | 1280x800 | 0.6383 | 3.3799 | 31.7335% | 0.2772% | 130 | 0 | 0 | none |
+| passage | 1280x800 | 1.6092 | 3.6878 | 50.4982% | 0.1733% | 52 | 0 | 0 | none |
+| far | 1280x800 | 0.7910 | 2.1578 | 44.8325% | 0.0198% | 128 | 0 | 0 | none |
+| passage_floor_closeup | 1280x800 | 0.9314 | 2.4800 | 43.0219% | 0.0500% | 100 | 0 | 0 | none |
+| tunnel_bend | 1280x800 | 0.8129 | 2.1599 | 40.9629% | 0.0708% | 55 | 0 | 0 | none |
+| remote_yaw0 | 1280x800 | 0.0416 | 1.0679 | 2.7991% | 0.0305% | 90 | 0 | 0 | none |
+| remote_yaw90 | 1280x800 | 0.0100 | 0.4323 | 0.9746% | 0.0035% | 94 | 0 | 0 | none |
+| remote_yaw180 | 1280x800 | 0.0470 | 1.3091 | 2.9463% | 0.0311% | 102 | 2 | 0 | (415,172)-(834,173) |
+| remote_yaw270 | 1280x800 | 0.0831 | 1.6885 | 5.1229% | 0.0580% | 107 | 2 | 0 | (835,204)-(886,205) |
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_final_r29_tilecontrol_comparison.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_final_r29_tilecontrol_comparison.md
new file mode 100644
index 0000000..f9d04b7
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_final_r29_tilecontrol_comparison.md
@@ -0,0 +1,9 @@
+# Step 17 LOD3 tile control comparison
+
+The `(6,2,0,lod3)` target was the only drawn tile, with the full desired clipmap retained. Captures were loaded as RGB arrays and compared numerically; full-size images were not opened in a viewer.
+
+`CPU rock to GPU background` counts pixels where the CPU target is rock `(96,86,68)` and the control is background `(14,19,27)`, the prior missing-coverage signature.
+
+| Control | Size | RGB MAE | RGB RMSE | Changed pixels | CPU rock to GPU background | Max channel |
+|---|---:|---:|---:|---:|---:|---:|
+| GPU final build r29 | 1280x800 | 0.0000 | 0.0000 | 0.0000% | 0 | 0 |
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_r26_tilecontrol_comparison.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_r26_tilecontrol_comparison.md
new file mode 100644
index 0000000..d39b8ce
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_r26_tilecontrol_comparison.md
@@ -0,0 +1,9 @@
+# Step 17 LOD3 tile control comparison
+
+The `(6,2,0,lod3)` target was the only drawn tile, with the full desired clipmap retained. Captures were loaded as RGB arrays and compared numerically; full-size images were not opened in a viewer.
+
+`CPU rock to GPU background` counts pixels where the CPU target is rock `(96,86,68)` and the control is background `(14,19,27)`, the prior missing-coverage signature.
+
+| Control | Size | RGB MAE | RGB RMSE | Changed pixels | CPU rock to GPU background | Max channel |
+|---|---:|---:|---:|---:|---:|---:|
+| GPU final build r26 | 1280x800 | 0.0000 | 0.0000 | 0.0000% | 0 | 0 |
diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_tilecontrol_comparison.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_tilecontrol_comparison.md
new file mode 100644
index 0000000..3836fad
--- /dev/null
+++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step17_tilecontrol_comparison.md
@@ -0,0 +1,18 @@
+# Step 17 LOD3 tile control comparison
+
+The `(6,2,0,lod3)` target was the only drawn tile, with the full desired clipmap retained. Captures were loaded as RGB arrays and compared numerically; full-size images were not opened in a viewer.
+
+`CPU rock to GPU background` counts pixels where the CPU target is rock `(96,86,68)` and the control is background `(14,19,27)`, the prior missing-coverage signature.
+
+| Control | Size | RGB MAE | RGB RMSE | Changed pixels | CPU rock to GPU background | Max channel |
+|---|---:|---:|---:|---:|---:|---:|
+| GPU indirect draw | 1280x800 | 0.9228 | 7.9137 | 1.4571% | 14921 | 82 |
+| GPU direct draw | 1280x800 | 0.9228 | 7.9137 | 1.4571% | 14921 | 82 |
+| GPU readback copy (cm) | 1280x800 | 0.9228 | 7.9137 | 1.4571% | 14921 | 82 |
+| GPU readback copy (m) | 1280x800 | 0.9228 | 7.9137 | 1.4571% | 14921 | 82 |
+| CPU mesh in GPU draw loop | 1280x800 | 0.0000 | 0.0000 | 0.0000% | 0 | 0 |
+| GPU depth test disabled | 1280x800 | 0.9325 | 7.9277 | 1.5908% | 14921 | 82 |
+| GPU skirts disabled | 1280x800 | 1.6339 | 10.5157 | 2.5859% | 26055 | 82 |
+| GPU connectors disabled | 1280x800 | 0.9228 | 7.9137 | 1.4571% | 14921 | 82 |
+| GPU skirt-depth fix | 1280x800 | 0.0000 | 0.0000 | 0.0000% | 0 | 0 |
+| GPU skirts and connectors disabled | 1280x800 | 1.6339 | 10.5157 | 2.5859% | 26055 | 82 |
diff --git a/Standalone/Source/StandaloneViewer.cpp b/Standalone/Source/StandaloneViewer.cpp
index 4b46cd5..0d82834 100644
--- a/Standalone/Source/StandaloneViewer.cpp
+++ b/Standalone/Source/StandaloneViewer.cpp
@@ -16,6 +16,7 @@
 #endif
 
 #include <d3d11.h>
+#include <d3d11_4.h>
 #include <d3dcompiler.h>
 #include <dxgi.h>
 #include <windowsx.h>
@@ -1362,6 +1363,9 @@ namespace
         ID3D11Buffer* IndirectArgs=nullptr;
         ID3D11Buffer* ConnectorVertexBuffer=nullptr;
         ID3D11Buffer* ConnectorIndexBuffer=nullptr;
+        ID3D11Buffer* ReadbackCopyVertexBuffer=nullptr;
+        ID3D11Buffer* ReadbackCopyIndexBuffer=nullptr;
+        uint32 ReadbackCopyIndexCount=0;
         ID3D11Query* CompletionEvent=nullptr;
         VoxelForgeGpuFarViewer::GpuTimingQueries TimingQueries{};
         uint64 VertexCapacityBytes=0;
@@ -1449,6 +1453,7 @@ namespace
         ID3D11Buffer* Constants=nullptr;
         ID3D11RasterizerState* SolidRasterizer=nullptr;
         ID3D11RasterizerState* WireRasterizer=nullptr;
+        ID3D11DepthStencilState* DepthDisabledAuditState=nullptr;
         std::shared_ptr<FStandaloneOwnerWorld> World;
         VoxelForgeCore::ResolvedWorldSnapshot Snapshot;
         std::unique_ptr<VoxelForgeCore::World> ShapeWorld;
@@ -1499,6 +1504,7 @@ namespace
         bool bShapeCompare=false;
         bool bFarShapeStagePending=false;
         bool bGpuFarEnabled=false;
+        bool bGpuFarPipelineInitialized=false;
         bool bDisableGpuColdStartDrain=false;
         bool bGpuFarAuditEnabled=false;
         bool bGpuFarShaderPrewarm=false,bDisableGpuWarmupPresentHold=false;
@@ -1512,7 +1518,10 @@ namespace
         std::atomic<bool> bGpuFarPreparationFinished{false};
         bool bGpuFarPreparationStarted=false,bGpuFarInputsFinalized=false;
         bool bGpuWarmupCompleted=false,bGpuFarTimelineEnabled=false;
+        bool bGpuWarmupRoutePending=false;
         bool bGpuShaderWarmupQueueSucceeded=false;
+        std::atomic<bool> bGpuShaderWarmupQueueFinished{false};
+        bool bGpuWarmupContextProtected=false;
         std::string GpuShaderWarmupQueueError;
         bool bGpuFarAuditReported=false;
         bool bGpuSkirtEmitEnabled=true,bGpuConnectorRepairEnabled=true;
@@ -1520,13 +1529,21 @@ namespace
         bool bForceBlockingStartupPresent=false;
         bool bPresentCpuFrameBeforeGpu=false;
         bool bCpuReadyFramePresentedBeforeGpu=false;
+        bool bQuitAfterClipmapClosePosted=false;
         bool bHoldCpuFrameDuringGpuWarmup=false,bGpuWarmupHoldReported=false;
-        bool bGpuWarmupReleaseReported=false;
+        bool bGpuWarmupReleaseReported=false,bGpuWarmupPresentSkipReported=false;
         uint64 GpuWarmupHeldPresents=0;
         bool bAuditGpuSkirtDrawArgs=false,bGpuSkirtDrawAuditReported=false;
+        bool bGpuFarDirectDrawAudit=false;
+        bool bGpuFarReadbackCopyDrawAudit=false;
+        bool bGpuFarMeterCopyDrawAudit=false;
+        bool bGpuFarCpuMeshDrawAudit=false;
+        bool bGpuFarDepthDisabledAudit=false,bGpuFarDepthDisabledAuditReported=false;
+        bool bDrawOnlyCpuTriangleAuditReported=false;
         ID3D11Buffer* GpuSkirtArgsStaging=nullptr;
         uint64 GpuSkirtDrawAuditTiles=0,GpuSkirtDrawAuditZeroCounts=0,
-            GpuSkirtDrawAuditOverflow=0,GpuSkirtDrawAuditBadStartIndex=0;
+            GpuSkirtDrawAuditOverflow=0,GpuSkirtDrawAuditBadStartIndex=0,
+            GpuSkirtDrawAuditBadSurfaceArgs=0;
         uint64 GpuSkirtDrawAuditSumIndices=0,GpuSkirtDrawAuditMaxIndices=0;
         std::unique_ptr<FFarSurvivalAudit> GpuFarAudit;
         bool bGpuFarFirstTileReported=false;
@@ -1547,7 +1564,9 @@ namespace
         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;
+        std::vector<FVoxelTileKey> GpuFarCandidateTiles,GpuFarBaseCpuFallbackTiles;
+        std::vector<FVoxelTileKey> GpuFarCpuFallbackTiles,GpuFarNeedPreparation;
+        uint64 GpuFarWarmupCpuClaimedTiles=0,GpuFarWarmupGpuTakeoverTiles=0;
         double GpuFarPreparationWallMs=0.0,GpuFarPreparationWorkerMs=0.0;
         uint64 GpuFarPumpCalls=0,GpuFarCompletionPolls=0;
         double GpuFarFirstPumpMs=-1.0,GpuFarLastPumpMs=-1.0;
@@ -1619,6 +1638,7 @@ namespace
             if(GpuShaderWarmupThread.joinable())GpuShaderWarmupThread.join();
             StopGeneration(true);ReleaseMeshes();
             if(SolidRasterizer)SolidRasterizer->Release();if(WireRasterizer)WireRasterizer->Release();
+            if(DepthDisabledAuditState)DepthDisabledAuditState->Release();
             if(Constants)Constants->Release();if(InputLayout)InputLayout->Release();
             if(GpuSkirtArgsStaging)GpuSkirtArgsStaging->Release();
             if(VertexShader)VertexShader->Release();if(PixelShader)PixelShader->Release();
@@ -1633,6 +1653,8 @@ namespace
             if(Tile.IndirectArgs)Tile.IndirectArgs->Release();
             if(Tile.ConnectorVertexBuffer)Tile.ConnectorVertexBuffer->Release();
             if(Tile.ConnectorIndexBuffer)Tile.ConnectorIndexBuffer->Release();
+            if(Tile.ReadbackCopyVertexBuffer)Tile.ReadbackCopyVertexBuffer->Release();
+            if(Tile.ReadbackCopyIndexBuffer)Tile.ReadbackCopyIndexBuffer->Release();
             if(Tile.CompletionEvent)Tile.CompletionEvent->Release();
             if(Tile.TimingQueries.start)Tile.TimingQueries.start->Release();
             if(Tile.TimingQueries.end)Tile.TimingQueries.end->Release();
@@ -1685,10 +1707,12 @@ namespace
             GpuFarBatches.clear();GpuFarQueue.clear();GpuFarQueued.clear();GpuFarEligibleTiles.clear();
             GpuFarFreeOutputSlots.clear();GpuFarNextOutputSlot=0;
             GpuFarPreparedInputs.clear();GpuFarPreparationErrors.clear();GpuFarPreparedAt.clear();
-            GpuFarCandidateTiles.clear();GpuFarCpuFallbackTiles.clear();GpuFarNeedPreparation.clear();
+            GpuFarCandidateTiles.clear();GpuFarBaseCpuFallbackTiles.clear();
+            GpuFarCpuFallbackTiles.clear();GpuFarNeedPreparation.clear();
             bGpuFarPreparationStarted=false;bGpuFarInputsFinalized=false;
             bGpuFarPreparationFinished.store(false,std::memory_order_relaxed);
-            bGpuWarmupCompleted=false;
+            bGpuWarmupRoutePending=false;
+            GpuFarWarmupCpuClaimedTiles=0;GpuFarWarmupGpuTakeoverTiles=0;
             GpuFarPreparationWallMs=0.0;GpuFarPreparationWorkerMs=0.0;
             GpuFarPumpCalls=0;GpuFarCompletionPolls=0;
             GpuFarFirstPumpMs=GpuFarLastPumpMs=GpuFarFirstSubmitMs=GpuFarLastSubmitMs=-1.0;
@@ -1701,6 +1725,9 @@ namespace
             GpuFarGpuDispatchMaxMs=0.0;GpuFarSubmissionSpanMs=0.0;GpuFarSubmissionSpanMaxMs=0.0;
             GpuFarLocalMemoryPeakBytes=0;GpuFarLocalMemoryBudgetBytes=0;
             bGpuFarFirstTileReported=false;
+            GpuSkirtDrawAuditTiles=GpuSkirtDrawAuditZeroCounts=GpuSkirtDrawAuditOverflow=0;
+            GpuSkirtDrawAuditBadStartIndex=GpuSkirtDrawAuditBadSurfaceArgs=0;
+            GpuSkirtDrawAuditSumIndices=GpuSkirtDrawAuditMaxIndices=0;
             GpuFarPoolReserveMs=0.0;
             GpuFarSubmissionBatches=0;GpuFarMaxTilesPerSubmissionBatch=0;
             DirtyBatchSet.clear();DesiredTiles.clear();DesiredTileSet.clear();CrossingWork.clear();
@@ -1896,49 +1923,79 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             D3D11_RASTERIZER_DESC Raster{};Raster.FillMode=D3D11_FILL_SOLID;Raster.CullMode=D3D11_CULL_NONE;
             Raster.DepthClipEnable=TRUE;if(FAILED(Device->CreateRasterizerState(&Raster,&SolidRasterizer)))return false;
             Raster.FillMode=D3D11_FILL_WIREFRAME;if(FAILED(Device->CreateRasterizerState(&Raster,&WireRasterizer)))return false;
+            D3D11_DEPTH_STENCIL_DESC NoDepth{};NoDepth.DepthEnable=FALSE;
+            NoDepth.DepthWriteMask=D3D11_DEPTH_WRITE_MASK_ZERO;NoDepth.DepthFunc=D3D11_COMPARISON_ALWAYS;
+            if(FAILED(Device->CreateDepthStencilState(&NoDepth,&DepthDisabledAuditState)))return false;
             if(!Resize(Width,Height))return false;
-            if(bGpuFarEnabled)
+            return true;
+        }
+
+        bool EnsureGpuFarPipeline(std::string& Error)
+        {
+            if(!bGpuFarEnabled||bGpuFarPipelineInitialized)return true;
+            if(!Device||!Context)
+            {Error="D3D device is unavailable for the GPU far pipeline";return false;}
+            const auto GpuPipelineStart=std::chrono::steady_clock::now();
+            std::printf("[ColdTimeline] stage=gpu_pipeline_init begin_ms=%.3f\n",
+                std::chrono::duration<double,std::milli>(GpuPipelineStart-ColdProcessStart).count());
+            if(!VoxelForgeGpuFarViewer::InitializeFarPipeline(Device,Context,GpuFarShaderPath,Error))
             {
-                std::string GpuInitError;
-                const auto GpuPipelineStart=std::chrono::steady_clock::now();
-                if(!VoxelForgeGpuFarViewer::InitializeFarPipeline(Device,Context,GpuFarShaderPath,
-                    GpuInitError))
-                {
-                    std::fprintf(stderr,"[GpuPipelineInitError] %s\n",GpuInitError.c_str());
-                    return false;
-                }
-                std::printf("[ColdTimeline] stage=gpu_pipeline_init end_ms=%.3f duration_ms=%.3f prewarm=%s\n",
-                    std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count(),
-                    std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-GpuPipelineStart).count(),
-                    bGpuFarShaderPrewarm?(bGpuFarWarmupResourcesOnly?"async_resources_only":"async_dispatch")
-                        :"disabled");
+                std::fprintf(stderr,"[GpuPipelineInitError] %s\n",Error.c_str());
+                return false;
             }
+            bGpuFarPipelineInitialized=true;
+            std::printf("[ColdTimeline] stage=gpu_pipeline_init end_ms=%.3f duration_ms=%.3f prewarm=%s\n",
+                std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count(),
+                std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-GpuPipelineStart).count(),
+                bGpuFarShaderPrewarm?(bGpuFarWarmupResourcesOnly?"async_resources_only":"async_dispatch")
+                    :"disabled");
+            return true;
+        }
+
+        bool SetGpuWarmupContextProtection(bool bProtected)
+        {
+            if(!Context)return false;
+            if(bGpuWarmupContextProtected==bProtected)return true;
+            ID3D11Multithread* Multithread=nullptr;
+            if(FAILED(Context->QueryInterface(__uuidof(ID3D11Multithread),
+                reinterpret_cast<void**>(&Multithread)))||!Multithread)return false;
+            Multithread->SetMultithreadProtected(bProtected?TRUE:FALSE);
+            Multithread->Release();
+            bGpuWarmupContextProtected=bProtected;
             return true;
         }
 
         void StartGpuShaderWarmupAsync()
         {
-            if(!bGpuFarShaderPrewarm||GpuShaderWarmupThread.joinable())return;
+            if(!bGpuFarShaderPrewarm||bGpuWarmupCompleted||GpuShaderWarmupThread.joinable())return;
             bGpuShaderWarmupQueueSucceeded=false;GpuShaderWarmupQueueError.clear();
+            bGpuShaderWarmupQueueFinished.store(false,std::memory_order_relaxed);
+            const bool bContextProtected=SetGpuWarmupContextProtection(true);
+            if(!bContextProtected)
+                std::fprintf(stderr,"[GpuWarmupContextProtection] enable_failed; async immediate-context protection is unavailable\n");
             GpuShaderWarmupThread=std::thread([this]()
             {
-                bool bNormalBasePriority=false;
-                if(bGpuFarWarmupNormalPriority)
-                {
-                    // The process stays BelowNormal; HIGHEST relative priority raises this
-                    // worker thread to a normal base priority on Windows.
-                    bNormalBasePriority=SetThreadPriority(GetCurrentThread(),THREAD_PRIORITY_HIGHEST)!=0;
-                    std::printf("[GpuShaderWarmupThread] priority=%s process=BelowNormal thread_relative=Highest\n",
-                        bNormalBasePriority?"Normal":"set_failed");
-                    std::fflush(stdout);
-                }
+                // Keep cold driver compilation from stealing CPU time from the CPU fallback
+                // workers. Cached shader warm-up is only a few milliseconds, so this does not
+                // delay the normal GPU route when the driver cache is warm.
+                const bool bPreserveAuditPriority=bGpuFarWarmupBeforeWorkers||bGpuFarAuditEnabled;
+                const int WarmupPriority=bGpuFarWarmupNormalPriority?THREAD_PRIORITY_HIGHEST:
+                    (bPreserveAuditPriority?THREAD_PRIORITY_BELOW_NORMAL:THREAD_PRIORITY_IDLE);
+                const bool bPrioritySet=SetThreadPriority(GetCurrentThread(),WarmupPriority)!=0;
+                std::printf("[GpuShaderWarmupThread] priority=%s process=BelowNormal thread_relative=%s\n",
+                    bPrioritySet?(bGpuFarWarmupNormalPriority?"Normal":
+                        (bPreserveAuditPriority?"BelowNormal":"Idle")):"set_failed",
+                    bGpuFarWarmupNormalPriority?"Highest":
+                        (bPreserveAuditPriority?"BelowNormal":"Idle"));
+                std::fflush(stdout);
                 bGpuShaderWarmupQueueSucceeded=VoxelForgeGpuFarViewer::QueueFarShaderWarmup(
                     Device,Context,GpuFarShaderPath,!bGpuFarWarmupResourcesOnly,GpuWarmupEntryFilter,
                     bGpuFarWarmupKernelFences,
                     GpuShaderWarmupQueueError);
+                bGpuShaderWarmupQueueFinished.store(true,std::memory_order_release);
             });
             std::printf("[GpuShaderWarmup] queue_thread=started overlap=%s\n",
-                bGpuFarWarmupBeforeWorkers?"none_before_generation_workers":"CPU_shape_compile");
+                bGpuFarWarmupBeforeWorkers?"none_before_generation_workers":"CPU_world_configuration");
             std::fflush(stdout);
         }
 
@@ -2195,6 +2252,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             StopGeneration();
             ReleaseMeshes();
             GenerationStart=std::chrono::steady_clock::now();
+            bQuitAfterClipmapClosePosted=false;
             std::printf("[ColdTimeline] stage=generation_setup begin_ms=%.3f\n",
                 std::chrono::duration<double,std::milli>(GenerationStart-ColdProcessStart).count());
             std::string Error;
@@ -2202,6 +2260,10 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             {std::fprintf(stderr,"World load failed: %s\n",Error.c_str());MessageBoxA(Window,Error.c_str(),"VoxelForge viewer",MB_ICONERROR);return;}
             std::printf("[ColdTimeline] stage=generation_setup configured_ms=%.3f\n",
                 std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count());
+            if(bGpuFarEnabled&&!EnsureGpuFarPipeline(Error))
+            {MessageBoxA(Window,Error.c_str(),"VoxelForge GPU pipeline",MB_ICONERROR);return;}
+            if(bGpuFarShaderPrewarm&&!bGpuFarWarmupBeforeWorkers)
+                StartGpuShaderWarmupAsync();
             const FVector CameraVoxels(Position.X*4.0,Position.Y*4.0,Position.Z*4.0);
             BuildTargetChunk=StandaloneWorldMesher::CameraChunk(CameraVoxels);
             bFullDetailReported=false;bInitialClipmapReported=false;bWholeViewPresented=false;bAutoCaptureWritten=false;
@@ -2580,8 +2642,12 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                                 Inputs[Index]->visibleSkirtFaceMask=World->bGenerateSkirts
                                     &&bGpuSkirtEmitEnabled
                                     ?ComputeVisibleSkirtFaceMask(Key,DesiredTileSet):0;
+                                // The CPU far mesher uses extent / CoarseTileCells, which can
+                                // differ from Key.StepVoxels() when coarse tiles use a reduced
+                                // cell count. Reuse the exported mesher step so skirt depth
+                                // stays identical to the CPU reference.
                                 Inputs[Index]->skirtDepthVoxels=std::max(1.0f,
-                                    static_cast<float>(World->SkirtCells))*Key.StepVoxels();
+                                    static_cast<float>(World->SkirtCells))*Inputs[Index]->step;
                                 if(!bGpuConnectorRepairEnabled)
                                 {
                                     Inputs[Index]->connectorPanels.clear();
@@ -2636,28 +2702,19 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         {
             GpuFarQueue.clear();GpuFarQueued.clear();GpuFarEligibleTiles.clear();
             GpuFarPreparedInputs.clear();GpuFarPreparationErrors.clear();GpuFarPreparedAt.clear();
-            GpuFarCandidateTiles.clear();GpuFarCpuFallbackTiles.clear();GpuFarNeedPreparation.clear();
+            GpuFarCandidateTiles.clear();GpuFarBaseCpuFallbackTiles.clear();
+            GpuFarCpuFallbackTiles.clear();GpuFarNeedPreparation.clear();
             bGpuFarPreparationStarted=false;bGpuFarInputsFinalized=false;
             bGpuFarPreparationFinished.store(false,std::memory_order_relaxed);
+            bGpuWarmupRoutePending=false;
             if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld())return;
             std::vector<FVoxelTileKey> CpuTiles;CpuTiles.reserve(Keys.size());
             std::vector<FVoxelTileKey> Eligible;Eligible.reserve(Keys.size());
             uint64 SkirtTransitionFallbackTiles=0;
-            uint64 RemoteLod3CpuFallbackTiles=0,RemoteLod4CpuFallbackTiles=0;
             for(const FVoxelTileKey& Key:Keys)
             {
                 const VoxelForgeCore::TileKey CoreKey{{Key.Coord.X,Key.Coord.Y,Key.Coord.Z},Key.Level};
                 const bool NearGpuEligible=Key.Level<World->FullResClipLevels;
-                if(Key.Level==3||Key.Level==4)
-                {
-                    CpuTiles.push_back(Key);
-                    if(Key.Level==3)++RemoteLod3CpuFallbackTiles;
-                    else ++RemoteLod4CpuFallbackTiles;
-                    if(bDrawOnlyTile&&IsDrawOnlyTile(Key))
-                        std::printf("[DrawOnlyTileRoute] key=(%d,%d,%d,lod%d) route=cpu_remote_band_fallback\n",
-                            Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level);
-                    continue;
-                }
                 if(NearGpuEligible||ShapeWorld->isFarFieldTileEligible(CoreKey))
                 {
                     const uint8 VisibleSkirtFaces=World->bGenerateSkirts
@@ -2698,10 +2755,6 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 ReleaseGpuFarTile(It->second);
                 It=GpuFarTiles.erase(It);
             }
-            if(RemoteLod3CpuFallbackTiles+RemoteLod4CpuFallbackTiles>0)
-                std::printf("[GpuRemoteBandFallback] lod3_cpu_tiles=%llu lod4_cpu_tiles=%llu reason=remote_capture_pixel_coverage\n",
-                    static_cast<unsigned long long>(RemoteLod3CpuFallbackTiles),
-                    static_cast<unsigned long long>(RemoteLod4CpuFallbackTiles));
             if(Eligible.size()>GpuFarOutputCapacitySlots)
             {
                 for(auto& Pair:GpuFarTiles)ReleaseGpuFarTile(Pair.second);
@@ -2719,7 +2772,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     CaptureOutputSlot,GpuFarOutputCapacitySlots);
             }
             GpuFarCandidateTiles=std::move(Eligible);
+            GpuFarBaseCpuFallbackTiles=CpuTiles;
             GpuFarCpuFallbackTiles=std::move(CpuTiles);
+            bGpuWarmupRoutePending=bGpuFarShaderPrewarm&&!bGpuFarAuditEnabled
+                &&!bGpuWarmupCompleted&&!bGpuFarWarmupBeforeWorkers
+                &&!GpuFarCandidateTiles.empty();
+            if(bGpuWarmupRoutePending)
+            {
+                GpuFarCpuFallbackTiles.reserve(GpuFarCpuFallbackTiles.size()+GpuFarCandidateTiles.size());
+                GpuFarCpuFallbackTiles.insert(GpuFarCpuFallbackTiles.end(),
+                    GpuFarCandidateTiles.begin(),GpuFarCandidateTiles.end());
+            }
             GpuFarNeedPreparation.reserve(GpuFarCandidateTiles.size());
             for(const FVoxelTileKey& Key:GpuFarCandidateTiles)
                 if(!GpuFarTiles.contains(Key))GpuFarNeedPreparation.push_back(Key);
@@ -2727,8 +2790,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             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("[GpuFieldEligibility] candidate_gpu_tiles=%zu near_gpu_tiles=%zu cpu_fallback_tiles=%zu warmup_cpu_fallback=%s\n",
+                GpuFarCandidateTiles.size(),NearGpuTiles,GpuFarCpuFallbackTiles.size(),
+                bGpuWarmupRoutePending?"yes":"no");
             std::printf("[GpuSkirtSeamFallback] near_tiles=%llu reason=fine_visible_lod_skirts gpu_far_skirts=active\n",
                 static_cast<unsigned long long>(SkirtTransitionFallbackTiles));
         }
@@ -2757,6 +2821,12 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         void StartGpuFarPreparation()
         {
             if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld()||bGpuFarPreparationStarted)return;
+            if(bGpuWarmupRoutePending&&!bGpuWarmupCompleted)
+            {
+                std::printf("[GpuFarPrepDeferred] reason=shader_warmup_cpu_tiles_own_candidates tiles=%zu\n",
+                    GpuFarCandidateTiles.size());
+                return;
+            }
             bGpuFarPreparationStarted=true;
             if(GpuFarCandidateTiles.empty())
             {
@@ -2829,33 +2899,77 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             GpuFarQueued.clear();for(const FVoxelTileKey& Key:GpuFarQueue)GpuFarQueued.insert(Key);
         }
 
+        bool CompleteGpuShaderWarmupAndRoute()
+        {
+            if(!bGpuFarShaderPrewarm||bGpuWarmupCompleted)return bGpuFarEnabled;
+            if(!bGpuShaderWarmupQueueFinished.load(std::memory_order_acquire))return false;
+            const auto WarmupStart=std::chrono::steady_clock::now();
+            if(GpuShaderWarmupThread.joinable())GpuShaderWarmupThread.join();
+            if(!bGpuShaderWarmupQueueSucceeded)
+            {
+                std::fprintf(stderr,"[GpuShaderWarmupError] queue: %s\n",
+                    GpuShaderWarmupQueueError.empty()?"warmup queue failed":GpuShaderWarmupQueueError.c_str());
+                bGpuFarEnabled=false;bGpuWarmupRoutePending=false;
+                SetGpuWarmupContextProtection(false);
+                AddGpuFarCpuFallbackTiles(DesiredTiles);
+                return false;
+            }
+            std::string WarmupError;
+            if(!VoxelForgeGpuFarViewer::CompleteFarShaderWarmup(Device,Context,
+                GpuFarShaderPath,WarmupError))
+            {
+                std::fprintf(stderr,"[GpuShaderWarmupError] %s\n",WarmupError.c_str());
+                bGpuFarEnabled=false;bGpuWarmupRoutePending=false;
+                SetGpuWarmupContextProtection(false);
+                AddGpuFarCpuFallbackTiles(DesiredTiles);
+                return false;
+            }
+            bGpuWarmupCompleted=true;
+            SetGpuWarmupContextProtection(false);
+            std::printf("[GpuShaderWarmupComplete] completion_wall_ms=%.3f\n",
+                std::chrono::duration<double,std::milli>(
+                    std::chrono::steady_clock::now()-WarmupStart).count());
+
+            if(!bGpuWarmupRoutePending||!Streamer)return true;
+            std::vector<FVoxelTileKey> ResidentKeys;ResidentKeys.reserve(ResidentTiles.size());
+            for(const auto& Pair:ResidentTiles)ResidentKeys.push_back(Pair.first);
+            std::vector<FVoxelTileKey> RetainedCandidates;
+            const FStandaloneStreamingUpdate Update=Streamer->RestrictDesiredTileFilterKeepingStarted(
+                GpuFarBaseCpuFallbackTiles,GpuFarCandidateTiles,ResidentKeys,RetainedCandidates);
+            for(const auto& Pair:Update.CancelledByCrossing)
+            {
+                auto WorkIt=CrossingWork.find(Pair.first);
+                if(WorkIt==CrossingWork.end())continue;
+                FCrossingWork& Work=WorkIt->second;
+                Work.PendingJobs=Work.PendingJobs>=Pair.second?Work.PendingJobs-Pair.second:0;
+                Work.CancelledQueuedJobs+=Pair.second;
+            }
+            std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> CpuOwned(
+                RetainedCandidates.begin(),RetainedCandidates.end());
+            const std::size_t CandidateCountBefore=GpuFarCandidateTiles.size();
+            GpuFarCandidateTiles.erase(std::remove_if(GpuFarCandidateTiles.begin(),GpuFarCandidateTiles.end(),
+                [&](const FVoxelTileKey& Key){return CpuOwned.contains(Key);}),GpuFarCandidateTiles.end());
+            GpuFarNeedPreparation.erase(std::remove_if(GpuFarNeedPreparation.begin(),GpuFarNeedPreparation.end(),
+                [&](const FVoxelTileKey& Key){return CpuOwned.contains(Key);}),GpuFarNeedPreparation.end());
+            GpuFarCpuFallbackTiles=GpuFarBaseCpuFallbackTiles;
+            GpuFarCpuFallbackTiles.insert(GpuFarCpuFallbackTiles.end(),
+                RetainedCandidates.begin(),RetainedCandidates.end());
+            GpuFarWarmupCpuClaimedTiles=RetainedCandidates.size();
+            GpuFarWarmupGpuTakeoverTiles=GpuFarCandidateTiles.size();
+            bGpuWarmupRoutePending=false;
+            const bool bHandoffValid=RetainedCandidates.size()+GpuFarCandidateTiles.size()==CandidateCountBefore;
+            std::printf("[GpuWarmupHandoff] cpu_started_ready_or_resident=%zu gpu_takeover_unstarted=%zu cancelled_cpu_queue_jobs=%d cpu_filter_tiles=%zu status=%s\n",
+                RetainedCandidates.size(),GpuFarCandidateTiles.size(),Update.CancelledQueuedJobs,
+                GpuFarCpuFallbackTiles.size(),bHandoffValid?"PASS":"FAIL");
+            return true;
+        }
+
         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();
@@ -2896,8 +3010,13 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
 
         void PumpGpuFarGeneration()
         {
-            if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld()||!Device||!Context
-                ||!FinalizeGpuFarPreparation()||GpuFarQueue.empty())return;
+            if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld()||!Device||!Context)return;
+            std::string PipelineError;
+            if(!EnsureGpuFarPipeline(PipelineError))
+            {bGpuFarEnabled=false;AddGpuFarCpuFallbackTiles(DesiredTiles);return;}
+            if(!CompleteGpuShaderWarmupAndRoute())return;
+            StartGpuFarPreparation();
+            if(!bGpuFarEnabled||!FinalizeGpuFarPreparation()||GpuFarQueue.empty())return;
             const bool bColdDrainBeforeWholeView=!bInitialClipmapReported&&!bGpuFarAuditEnabled
                 &&!bDisableGpuColdStartDrain;
             const auto PumpStart=std::chrono::steady_clock::now();
@@ -2909,7 +3028,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 GpuFarPumpGapMs+=GapMs;GpuFarPumpGapMaxMs=std::max(GpuFarPumpGapMaxMs,GapMs);
             }
             LastGpuFarPumpAt=PumpStart;GpuFarLastPumpMs=PumpMs;++GpuFarPumpCalls;
-            const bool bDeferredBatch=!bGpuFarAuditEnabled;
+            const bool bDeferredBatch=!bGpuFarAuditEnabled
+                &&!(bAuditGpuSkirtDrawArgs&&bDrawOnlyTile);
             if(bDeferredBatch)
             {
                 std::string BatchError;
@@ -2938,6 +3058,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     bGpuFarEnabled=false;break;
                 }
                 VoxelForgeGpuFarViewer::CpuMesh AuditMesh;
+                const bool bAuditDrawOnlyGpuTile=bAuditGpuSkirtDrawArgs&&bDrawOnlyTile
+                    &&IsDrawOnlyTile(Key);
                 std::string Error;double EnqueueMs=0.0,GpuDispatchMs=0.0;
                 const VoxelForgeCore::TileKey CoreKey{{Key.Coord.X,Key.Coord.Y,Key.Coord.Z},Key.Level};
                 const auto PreparedInputIt=GpuFarPreparedInputs.find(Key);
@@ -2961,7 +3083,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     Tile.OutputSlot,GpuFarOutputCapacitySlots,&Tile.IndirectArgsByteOffset,
                     bDeferredBatch?nullptr:&Tile.CompletionEvent,
                     &Tile.TimingQueries,
-                    (bGpuFarAuditEnabled&&Key.Level>=6)?&AuditMesh:nullptr,
+                    ((bGpuFarAuditEnabled&&Key.Level>=6)||bAuditDrawOnlyGpuTile)?&AuditMesh:nullptr,
                     &EnqueueMs,&GpuDispatchMs,Error))
                 {
                     GpuFarFreeOutputSlots.push_back(Tile.OutputSlot);
@@ -2971,6 +3093,405 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     bGpuFarEnabled=false;
                     break;
                 }
+                if(bAuditDrawOnlyGpuTile)
+                    std::printf("[GpuPooledOutputReadback] key=(%d,%d,%d,lod%d) output_slot=%u indirect_offset=%u surface_indices=%u output_vertices=%zu output_indices=%zu sequential_indices=%d indirect_args_valid=%d vertex_stride_bytes=%u status=%s\n",
+                        Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,Tile.OutputSlot,
+                        Tile.IndirectArgsByteOffset,AuditMesh.surfaceVertexCount,
+                        AuditMesh.vertices.size(),AuditMesh.indices.size(),
+                        AuditMesh.sequentialIndices,AuditMesh.indirectArgsValid,
+                        static_cast<unsigned>(sizeof(FGpuVertex)),
+                        AuditMesh.vertices.size()==AuditMesh.indices.size()
+                            &&AuditMesh.sequentialIndices&&AuditMesh.indirectArgsValid?"PASS":"FAIL");
+                if(bAuditDrawOnlyGpuTile)
+                {
+                    std::string CpuTriangleError;
+                    const FVector TriangleAuditCameraVoxels(Position.X*4.0,Position.Y*4.0,Position.Z*4.0);
+                    const FStandaloneDesiredTileSet TriangleAuditDesired=
+                        StandaloneWorldMesher::BuildDesiredSet(*World,TriangleAuditCameraVoxels);
+                    int32 CpuBandMin=MIN_int32,CpuBandMax=MAX_int32;
+                    if(World->StrateContentCutMinLevel<9&&TriangleAuditDesired.bHasStrate)
+                    {
+                        CpuBandMin=TriangleAuditDesired.StrateBottomChunkZ*CHUNK_SIZE;
+                        CpuBandMax=(TriangleAuditDesired.StrateTopChunkZ+1)*CHUNK_SIZE;
+                    }
+                    VoxelForgeCore::TileGenerationOptions CpuTriangleOptions;
+                    CpuTriangleOptions.bandMinZVoxel=CpuBandMin;
+                    CpuTriangleOptions.bandMaxZVoxel=CpuBandMax;
+                    auto CpuTriangleMesh=ShapeWorld->generateTile(CoreKey,
+                        CpuTriangleOptions,CpuTriangleError);
+                    if(!CpuTriangleMesh)
+                    {
+                        std::printf("[GpuCpuTriangleAudit] key=(%d,%d,%d,lod%d) status=FAIL error=%s\n",
+                            Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,CpuTriangleError.c_str());
+                    }
+                    else
+                    {
+                        using FPoint=std::array<double,3>;
+                        using FTriangle=std::array<FPoint,3>;
+                        std::vector<FTriangle> CpuTriangles,GpuTriangles,CpuTriangleNormals,
+                            GpuTriangleNormals,CpuSkirtTriangles,GpuSkirtTriangles;
+                        std::size_t CpuTriangleCount=0,CpuVisibleSkirtTriangleCount=0;
+                        const std::size_t GpuSurfaceVertexCount=std::min<std::size_t>(
+                            AuditMesh.surfaceVertexCount,AuditMesh.vertices.size());
+                        const std::size_t GpuSkirtStart=std::min<std::size_t>(
+                            GpuSurfaceVertexCount+AuditMesh.connectorPanelVertexCount,
+                            AuditMesh.vertices.size());
+                        const std::size_t GpuSkirtEnd=std::min<std::size_t>(
+                            GpuSkirtStart+AuditMesh.skirtVertexCount,AuditMesh.vertices.size());
+                        const std::size_t GpuTriangleCount=GpuSurfaceVertexCount/3;
+                        CpuTriangles.reserve(CpuTriangleMesh->indices.size()/3);
+                        GpuTriangles.reserve(GpuTriangleCount);
+                        CpuSkirtTriangles.reserve(CpuTriangleMesh->groups.skirtTriangleCount);
+                        GpuSkirtTriangles.reserve((GpuSkirtEnd-GpuSkirtStart)/3);
+                        for(std::size_t I=0;I+2<CpuTriangleMesh->indices.size();I+=3)
+                        {
+                            const std::size_t TriangleIndex=I/3;
+                            const uint8 Faces=TriangleIndex<CpuTriangleMesh->skirtFaceMasks.size()
+                                ?CpuTriangleMesh->skirtFaceMasks[TriangleIndex]:0;
+                            const auto IA=CpuTriangleMesh->indices[I];
+                            const auto IB=CpuTriangleMesh->indices[I+1];
+                            const auto IC=CpuTriangleMesh->indices[I+2];
+                            if(IA>=CpuTriangleMesh->vertices.size()||IB>=CpuTriangleMesh->vertices.size()
+                                ||IC>=CpuTriangleMesh->vertices.size())continue;
+                            const auto& A=CpuTriangleMesh->vertices[IA].position;
+                            const auto& B=CpuTriangleMesh->vertices[IB].position;
+                            const auto& C=CpuTriangleMesh->vertices[IC].position;
+                            if(Faces==0)
+                            {
+                                CpuTriangles.push_back({FPoint{A.x,A.y,A.z},FPoint{B.x,B.y,B.z},
+                                    FPoint{C.x,C.y,C.z}});
+                                const auto& NA=CpuTriangleMesh->vertices[IA].normal;
+                                const auto& NB=CpuTriangleMesh->vertices[IB].normal;
+                                const auto& NC=CpuTriangleMesh->vertices[IC].normal;
+                                CpuTriangleNormals.push_back({FPoint{NA.x,NA.y,NA.z},
+                                    FPoint{NB.x,NB.y,NB.z},FPoint{NC.x,NC.y,NC.z}});
+                                ++CpuTriangleCount;
+                            }
+                            else if((Faces&PreparedSkirtFaceMask)!=0)
+                            {
+                                CpuSkirtTriangles.push_back({FPoint{A.x,A.y,A.z},FPoint{B.x,B.y,B.z},
+                                    FPoint{C.x,C.y,C.z}});
+                                ++CpuVisibleSkirtTriangleCount;
+                            }
+                        }
+                        for(std::size_t I=0;I+2<GpuSurfaceVertexCount;I+=3)
+                        {
+                            const auto& A=AuditMesh.vertices[I].position;
+                            const auto& B=AuditMesh.vertices[I+1].position;
+                            const auto& C=AuditMesh.vertices[I+2].position;
+                            GpuTriangles.push_back({FPoint{A[0],A[1],A[2]},FPoint{B[0],B[1],B[2]},
+                                FPoint{C[0],C[1],C[2]}});
+                            const auto& NA=AuditMesh.vertices[I].normal;
+                            const auto& NB=AuditMesh.vertices[I+1].normal;
+                            const auto& NC=AuditMesh.vertices[I+2].normal;
+                            GpuTriangleNormals.push_back({FPoint{NA[0],NA[1],NA[2]},
+                                FPoint{NB[0],NB[1],NB[2]},FPoint{NC[0],NC[1],NC[2]}});
+                        }
+                        for(std::size_t I=GpuSkirtStart;I+2<GpuSkirtEnd;I+=3)
+                        {
+                            const auto& A=AuditMesh.vertices[I].position;
+                            const auto& B=AuditMesh.vertices[I+1].position;
+                            const auto& C=AuditMesh.vertices[I+2].position;
+                            GpuSkirtTriangles.push_back({FPoint{A[0],A[1],A[2]},FPoint{B[0],B[1],B[2]},
+                                FPoint{C[0],C[1],C[2]}});
+                        }
+                        const double Infinity=std::numeric_limits<double>::infinity();
+                        FPoint CpuSkirtMinimum{Infinity,Infinity,Infinity};
+                        FPoint CpuSkirtMaximum{-Infinity,-Infinity,-Infinity};
+                        FPoint GpuSkirtMinimum{Infinity,Infinity,Infinity};
+                        FPoint GpuSkirtMaximum{-Infinity,-Infinity,-Infinity};
+                        const auto AccumulateBounds=[](const std::vector<FTriangle>& Triangles,
+                            FPoint& Minimum,FPoint& Maximum)
+                        {
+                            for(const FTriangle& Triangle:Triangles)
+                                for(const FPoint& Point:Triangle)
+                                    for(int Axis=0;Axis<3;++Axis)
+                                    {
+                                        Minimum[Axis]=std::min(Minimum[Axis],Point[Axis]);
+                                        Maximum[Axis]=std::max(Maximum[Axis],Point[Axis]);
+                                    }
+                        };
+                        AccumulateBounds(CpuSkirtTriangles,CpuSkirtMinimum,CpuSkirtMaximum);
+                        AccumulateBounds(GpuSkirtTriangles,GpuSkirtMinimum,GpuSkirtMaximum);
+                        constexpr double TriangleVertexToleranceCm=0.05;
+                        constexpr double TriangleVertexToleranceSq=
+                            TriangleVertexToleranceCm*TriangleVertexToleranceCm;
+                        constexpr int TrianglePermutations[6][3]={{0,1,2},{0,2,1},{1,0,2},
+                            {1,2,0},{2,0,1},{2,1,0}};
+                        std::vector<uint8> MatchedGpuTriangles(GpuTriangles.size(),0);
+                        int64 Missing=0,Extra=0;
+                        double MaxMatchedVertexErrorCm=0.0;
+                        double MaxMatchedSurfaceNormalDelta=0.0,MinMatchedSurfaceNormalDot=1.0;
+                        for(std::size_t CpuIndex=0;CpuIndex<CpuTriangles.size();++CpuIndex)
+                        {
+                            const FTriangle& CpuTriangle=CpuTriangles[CpuIndex];
+                            double BestCost=std::numeric_limits<double>::infinity();
+                            double BestMaxDistanceSq=0.0;
+                            std::size_t BestGpuIndex=GpuTriangles.size();
+                            int BestPermutation[3]={0,1,2};
+                            for(std::size_t GpuIndex=0;GpuIndex<GpuTriangles.size();++GpuIndex)
+                            {
+                                if(MatchedGpuTriangles[GpuIndex])continue;
+                                const FTriangle& GpuTriangle=GpuTriangles[GpuIndex];
+                                for(const auto& Permutation:TrianglePermutations)
+                                {
+                                    double Cost=0.0,MaxDistanceSq=0.0;
+                                    for(int Vertex=0;Vertex<3;++Vertex)
+                                    {
+                                        double DistanceSq=0.0;
+                                        for(int Axis=0;Axis<3;++Axis)
+                                        {
+                                            const double Delta=CpuTriangle[Vertex][Axis]
+                                                -GpuTriangle[Permutation[Vertex]][Axis];
+                                            DistanceSq+=Delta*Delta;
+                                        }
+                                        Cost+=DistanceSq;MaxDistanceSq=std::max(MaxDistanceSq,DistanceSq);
+                                    }
+                                    if(MaxDistanceSq<=TriangleVertexToleranceSq&&Cost<BestCost)
+                                    {
+                                        BestCost=Cost;BestMaxDistanceSq=MaxDistanceSq;
+                                        BestGpuIndex=GpuIndex;
+                                        for(int Vertex=0;Vertex<3;++Vertex)
+                                            BestPermutation[Vertex]=Permutation[Vertex];
+                                    }
+                                }
+                            }
+                            if(BestGpuIndex==GpuTriangles.size())++Missing;
+                            else
+                            {
+                                MatchedGpuTriangles[BestGpuIndex]=1;
+                                MaxMatchedVertexErrorCm=std::max(MaxMatchedVertexErrorCm,
+                                    std::sqrt(BestMaxDistanceSq));
+                                for(int Vertex=0;Vertex<3;++Vertex)
+                                {
+                                    const FPoint& CpuNormal=CpuTriangleNormals[CpuIndex][Vertex];
+                                    const FPoint& GpuNormal=GpuTriangleNormals[BestGpuIndex]
+                                        [BestPermutation[Vertex]];
+                                    double DeltaSq=0.0,CpuLengthSq=0.0,GpuLengthSq=0.0,Dot=0.0;
+                                    for(int Axis=0;Axis<3;++Axis)
+                                    {
+                                        const double Delta=CpuNormal[Axis]-GpuNormal[Axis];
+                                        DeltaSq+=Delta*Delta;
+                                        CpuLengthSq+=CpuNormal[Axis]*CpuNormal[Axis];
+                                        GpuLengthSq+=GpuNormal[Axis]*GpuNormal[Axis];
+                                        Dot+=CpuNormal[Axis]*GpuNormal[Axis];
+                                    }
+                                    MaxMatchedSurfaceNormalDelta=std::max(MaxMatchedSurfaceNormalDelta,
+                                        std::sqrt(DeltaSq));
+                                    if(CpuLengthSq>1.0e-20&&GpuLengthSq>1.0e-20)
+                                        MinMatchedSurfaceNormalDot=std::min(MinMatchedSurfaceNormalDot,
+                                            Dot/std::sqrt(CpuLengthSq*GpuLengthSq));
+                                }
+                            }
+                        }
+                        for(const uint8 Matched:MatchedGpuTriangles)if(!Matched)++Extra;
+                        std::vector<uint8> MatchedGpuSkirts(GpuSkirtTriangles.size(),0);
+                        int64 MissingSkirts=0,ExtraSkirts=0;
+                        double MaxMatchedSkirtVertexErrorCm=0.0;
+                        double MinNearestUnmatchedSkirtErrorCm=Infinity;
+                        double MaxNearestUnmatchedSkirtErrorCm=0.0;
+                        for(const FTriangle& CpuTriangle:CpuSkirtTriangles)
+                        {
+                            double BestCost=std::numeric_limits<double>::infinity();
+                            double BestMaxDistanceSq=0.0;
+                            double NearestMaxDistanceSq=std::numeric_limits<double>::infinity();
+                            std::size_t BestGpuIndex=GpuSkirtTriangles.size();
+                            for(std::size_t GpuIndex=0;GpuIndex<GpuSkirtTriangles.size();++GpuIndex)
+                            {
+                                if(MatchedGpuSkirts[GpuIndex])continue;
+                                const FTriangle& GpuTriangle=GpuSkirtTriangles[GpuIndex];
+                                for(const auto& Permutation:TrianglePermutations)
+                                {
+                                    double Cost=0.0,MaxDistanceSq=0.0;
+                                    for(int Vertex=0;Vertex<3;++Vertex)
+                                    {
+                                        double DistanceSq=0.0;
+                                        for(int Axis=0;Axis<3;++Axis)
+                                        {
+                                            const double Delta=CpuTriangle[Vertex][Axis]
+                                                -GpuTriangle[Permutation[Vertex]][Axis];
+                                            DistanceSq+=Delta*Delta;
+                                        }
+                                        Cost+=DistanceSq;MaxDistanceSq=std::max(MaxDistanceSq,DistanceSq);
+                                    }
+                                    NearestMaxDistanceSq=std::min(NearestMaxDistanceSq,MaxDistanceSq);
+                                    if(MaxDistanceSq<=TriangleVertexToleranceSq&&Cost<BestCost)
+                                    {BestCost=Cost;BestMaxDistanceSq=MaxDistanceSq;BestGpuIndex=GpuIndex;}
+                                }
+                            }
+                            if(BestGpuIndex==GpuSkirtTriangles.size())
+                            {
+                                ++MissingSkirts;
+                                if(std::isfinite(NearestMaxDistanceSq))
+                                {
+                                    const double NearestErrorCm=std::sqrt(NearestMaxDistanceSq);
+                                    MinNearestUnmatchedSkirtErrorCm=std::min(
+                                        MinNearestUnmatchedSkirtErrorCm,NearestErrorCm);
+                                    MaxNearestUnmatchedSkirtErrorCm=std::max(
+                                        MaxNearestUnmatchedSkirtErrorCm,NearestErrorCm);
+                                }
+                            }
+                            else
+                            {
+                                MatchedGpuSkirts[BestGpuIndex]=1;
+                                MaxMatchedSkirtVertexErrorCm=std::max(MaxMatchedSkirtVertexErrorCm,
+                                    std::sqrt(BestMaxDistanceSq));
+                            }
+                        }
+                        for(const uint8 Matched:MatchedGpuSkirts)if(!Matched)++ExtraSkirts;
+                        if(!CpuSkirtTriangles.empty()&&!GpuSkirtTriangles.empty())
+                        {
+                            const FTriangle& CpuSample=CpuSkirtTriangles.front();
+                            std::size_t NearestGpu=GpuSkirtTriangles.size();
+                            int NearestPermutation[3]={0,1,2};
+                            double NearestMaxDistanceSq=Infinity;
+                            for(std::size_t GpuIndex=0;GpuIndex<GpuSkirtTriangles.size();++GpuIndex)
+                                for(const auto& Permutation:TrianglePermutations)
+                                {
+                                    double MaxDistanceSq=0.0;
+                                    for(int Vertex=0;Vertex<3;++Vertex)
+                                    {
+                                        double DistanceSq=0.0;
+                                        for(int Axis=0;Axis<3;++Axis)
+                                        {
+                                            const double Delta=CpuSample[Vertex][Axis]
+                                                -GpuSkirtTriangles[GpuIndex][Permutation[Vertex]][Axis];
+                                            DistanceSq+=Delta*Delta;
+                                        }
+                                        MaxDistanceSq=std::max(MaxDistanceSq,DistanceSq);
+                                    }
+                                    if(MaxDistanceSq<NearestMaxDistanceSq)
+                                    {
+                                        NearestMaxDistanceSq=MaxDistanceSq;NearestGpu=GpuIndex;
+                                        for(int Vertex=0;Vertex<3;++Vertex)
+                                            NearestPermutation[Vertex]=Permutation[Vertex];
+                                    }
+                                }
+                            if(NearestGpu<GpuSkirtTriangles.size())
+                                std::printf("[GpuCpuTriangleAuditSkirtSample] cpu=(%.3f,%.3f,%.3f)|(%.3f,%.3f,%.3f)|(%.3f,%.3f,%.3f) nearest_gpu=%zu gpu=(%.3f,%.3f,%.3f)|(%.3f,%.3f,%.3f)|(%.3f,%.3f,%.3f) max_error_cm=%.6f\n",
+                                    CpuSample[0][0],CpuSample[0][1],CpuSample[0][2],
+                                    CpuSample[1][0],CpuSample[1][1],CpuSample[1][2],
+                                    CpuSample[2][0],CpuSample[2][1],CpuSample[2][2],NearestGpu,
+                                    GpuSkirtTriangles[NearestGpu][NearestPermutation[0]][0],
+                                    GpuSkirtTriangles[NearestGpu][NearestPermutation[0]][1],
+                                    GpuSkirtTriangles[NearestGpu][NearestPermutation[0]][2],
+                                    GpuSkirtTriangles[NearestGpu][NearestPermutation[1]][0],
+                                    GpuSkirtTriangles[NearestGpu][NearestPermutation[1]][1],
+                                    GpuSkirtTriangles[NearestGpu][NearestPermutation[1]][2],
+                                    GpuSkirtTriangles[NearestGpu][NearestPermutation[2]][0],
+                                    GpuSkirtTriangles[NearestGpu][NearestPermutation[2]][1],
+                                    GpuSkirtTriangles[NearestGpu][NearestPermutation[2]][2],
+                                    std::sqrt(NearestMaxDistanceSq));
+                        }
+                        std::printf("[GpuCpuTriangleAudit] key=(%d,%d,%d,lod%d) cpu_band_min_z_voxel=%d cpu_band_max_z_voxel=%d cpu_surface_triangles=%zu gpu_surface_triangles=%zu cpu_visible_skirt_triangles=%zu gpu_skirt_triangles=%zu cpu_all_skirt_triangles=%u visible_skirt_mask=0x%02x unmatched_cpu=%lld unmatched_gpu=%lld unmatched_cpu_skirts=%lld unmatched_gpu_skirts=%lld tolerance_cm=%.3f max_surface_vertex_error_cm=%.8g max_matched_surface_normal_delta=%.8g min_matched_surface_normal_dot=%.8g max_skirt_vertex_error_cm=%.8g nearest_unmatched_skirt_error_cm_min=%.8g nearest_unmatched_skirt_error_cm_max=%.8g cpu_skirt_bounds_min_cm=(%.6f,%.6f,%.6f) cpu_skirt_bounds_max_cm=(%.6f,%.6f,%.6f) gpu_skirt_bounds_min_cm=(%.6f,%.6f,%.6f) gpu_skirt_bounds_max_cm=(%.6f,%.6f,%.6f) status=%s\n",
+                            Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,CpuBandMin,CpuBandMax,
+                            CpuTriangleCount,GpuTriangleCount,CpuVisibleSkirtTriangleCount,
+                            GpuSkirtTriangles.size(),
+                            CpuTriangleMesh->groups.skirtTriangleCount,PreparedSkirtFaceMask,
+                            static_cast<long long>(Missing),static_cast<long long>(Extra),
+                            static_cast<long long>(MissingSkirts),static_cast<long long>(ExtraSkirts),
+                            TriangleVertexToleranceCm,MaxMatchedVertexErrorCm,
+                            MaxMatchedSurfaceNormalDelta,MinMatchedSurfaceNormalDot,
+                            MaxMatchedSkirtVertexErrorCm,
+                            std::isfinite(MinNearestUnmatchedSkirtErrorCm)
+                                ?MinNearestUnmatchedSkirtErrorCm:0.0,
+                            MaxNearestUnmatchedSkirtErrorCm,
+                            CpuSkirtTriangles.empty()?0.0:CpuSkirtMinimum[0],
+                            CpuSkirtTriangles.empty()?0.0:CpuSkirtMinimum[1],
+                            CpuSkirtTriangles.empty()?0.0:CpuSkirtMinimum[2],
+                            CpuSkirtTriangles.empty()?0.0:CpuSkirtMaximum[0],
+                            CpuSkirtTriangles.empty()?0.0:CpuSkirtMaximum[1],
+                            CpuSkirtTriangles.empty()?0.0:CpuSkirtMaximum[2],
+                            GpuSkirtTriangles.empty()?0.0:GpuSkirtMinimum[0],
+                            GpuSkirtTriangles.empty()?0.0:GpuSkirtMinimum[1],
+                            GpuSkirtTriangles.empty()?0.0:GpuSkirtMinimum[2],
+                            GpuSkirtTriangles.empty()?0.0:GpuSkirtMaximum[0],
+                            GpuSkirtTriangles.empty()?0.0:GpuSkirtMaximum[1],
+                            GpuSkirtTriangles.empty()?0.0:GpuSkirtMaximum[2],
+                            Missing==0&&Extra==0&&MissingSkirts==0&&ExtraSkirts==0
+                                &&CpuTriangleCount==GpuTriangleCount
+                                &&CpuVisibleSkirtTriangleCount==GpuSkirtTriangles.size()?"PASS":"FAIL");
+                        if(bGpuFarCpuMeshDrawAudit&&bDrawOnlyTile&&IsDrawOnlyTile(Key))
+                        {
+                            std::vector<FGpuVertex> CpuDrawVertices;CpuDrawVertices.reserve(
+                                CpuTriangleMesh->vertices.size());
+                            for(const auto& Vertex:CpuTriangleMesh->vertices)
+                            {
+                                FGpuVertex Copy{};
+                                Copy.Position[0]=static_cast<float>(Vertex.position.x*0.01);
+                                Copy.Position[1]=static_cast<float>(Vertex.position.y*0.01);
+                                Copy.Position[2]=static_cast<float>(Vertex.position.z*0.01);
+                                Copy.Normal[0]=static_cast<float>(Vertex.normal.x);
+                                Copy.Normal[1]=static_cast<float>(Vertex.normal.y);
+                                Copy.Normal[2]=static_cast<float>(Vertex.normal.z);
+                                CpuDrawVertices.push_back(Copy);
+                            }
+                            std::vector<uint32> CpuDrawIndices;CpuDrawIndices.reserve(
+                                CpuTriangleMesh->indices.size());
+                            for(size_t I=0;I+2<CpuTriangleMesh->indices.size();I+=3)
+                            {
+                                const size_t TriangleIndex=I/3;
+                                const uint8 Faces=TriangleIndex<CpuTriangleMesh->skirtFaceMasks.size()
+                                    ?CpuTriangleMesh->skirtFaceMasks[TriangleIndex]:0;
+                                if(Faces!=0&&(Faces&PreparedSkirtFaceMask)==0)continue;
+                                CpuDrawIndices.push_back(CpuTriangleMesh->indices[I]);
+                                CpuDrawIndices.push_back(CpuTriangleMesh->indices[I+1]);
+                                CpuDrawIndices.push_back(CpuTriangleMesh->indices[I+2]);
+                            }
+                            if(!MakeBuffer(&Tile.ReadbackCopyVertexBuffer,CpuDrawVertices.data(),
+                                CpuDrawVertices.size()*sizeof(FGpuVertex),D3D11_BIND_VERTEX_BUFFER)
+                                ||!MakeBuffer(&Tile.ReadbackCopyIndexBuffer,CpuDrawIndices.data(),
+                                    CpuDrawIndices.size()*sizeof(uint32),D3D11_BIND_INDEX_BUFFER))
+                            {
+                                std::fprintf(stderr,"[GpuCpuMeshDrawAudit] key=(%d,%d,%d,lod%d) buffer_creation=FAIL\n",
+                                    Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level);
+                                bGpuFarEnabled=false;
+                            }
+                            else
+                            {
+                                Tile.ReadbackCopyIndexCount=static_cast<uint32>(CpuDrawIndices.size());
+                                std::printf("[GpuCpuMeshDrawAudit] key=(%d,%d,%d,lod%d) surface_triangles=%zu all_skirt_triangles=%u visible_skirt_mask=0x%02x draw_indices=%u vertices=%zu path=cpu_mesh_copy\n",
+                                    Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,CpuTriangleCount,
+                                    CpuTriangleMesh->groups.skirtTriangleCount,PreparedSkirtFaceMask,
+                                    Tile.ReadbackCopyIndexCount,CpuDrawVertices.size());
+                            }
+                        }
+                    }
+                }
+                if((bGpuFarReadbackCopyDrawAudit||bGpuFarMeterCopyDrawAudit)
+                    &&bDrawOnlyTile&&IsDrawOnlyTile(Key)
+                    &&!AuditMesh.vertices.empty()&&AuditMesh.vertices.size()==AuditMesh.indices.size())
+                {
+                    std::vector<FGpuVertex> CopyVertices;CopyVertices.reserve(AuditMesh.vertices.size());
+                    for(const auto& Vertex:AuditMesh.vertices)
+                    {
+                        FGpuVertex Copy{};
+                        for(int Axis=0;Axis<3;++Axis)
+                        {
+                            Copy.Position[Axis]=Vertex.position[Axis]
+                                *(bGpuFarMeterCopyDrawAudit?0.01f:1.0f);
+                            Copy.Normal[Axis]=Vertex.normal[Axis];
+                        }
+                        CopyVertices.push_back(Copy);
+                    }
+                    if(!MakeBuffer(&Tile.ReadbackCopyVertexBuffer,CopyVertices.data(),
+                        CopyVertices.size()*sizeof(FGpuVertex),D3D11_BIND_VERTEX_BUFFER)
+                        ||!MakeBuffer(&Tile.ReadbackCopyIndexBuffer,AuditMesh.indices.data(),
+                            AuditMesh.indices.size()*sizeof(uint32),D3D11_BIND_INDEX_BUFFER))
+                    {
+                        std::fprintf(stderr,"[GpuReadbackCopyDrawAudit] key=(%d,%d,%d,lod%d) buffer_creation=FAIL\n",
+                            Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level);
+                        bGpuFarEnabled=false;
+                    }
+                    else
+                    {
+                        Tile.ReadbackCopyIndexCount=static_cast<uint32>(AuditMesh.indices.size());
+                        std::printf("[GpuReadbackCopyDrawAudit] key=(%d,%d,%d,lod%d) indices=%u vertices=%zu coordinates=%s path=unpooled_readback_copy\n",
+                            Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,
+                            Tile.ReadbackCopyIndexCount,CopyVertices.size(),
+                            bGpuFarMeterCopyDrawAudit?"meters":"centimeters");
+                    }
+                }
                 if(PreparedInputIt!=GpuFarPreparedInputs.end())GpuFarPreparedInputs.erase(PreparedInputIt);
                 if(bGpuFarAuditEnabled&&GpuFarAudit&&Key.Level>=6)
                 {
@@ -3093,6 +3614,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             bGpuFarEnabled=!bGpuFarEnabled;
             if(bGpuFarEnabled)
             {
+                std::string GpuInitError;
+                if(!EnsureGpuFarPipeline(GpuInitError))
+                {bGpuFarEnabled=false;return;}
                 bPresentCpuFrameBeforeGpu=true;
                 bHoldCpuFrameDuringGpuWarmup=true;
                 bCpuReadyFramePresentedBeforeGpu=false;
@@ -3427,6 +3951,30 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     Indices.push_back(Base+Resident.Indices[Start+1]);
                     Indices.push_back(Base+Resident.Indices[Start+2]);
                 }
+                if(bAuditGpuSkirtDrawArgs&&bDrawOnlyTile&&IsDrawOnlyTile(Tile)
+                    &&!bDrawOnlyCpuTriangleAuditReported)
+                {
+                    uint64 SurfaceTriangles=0,AllSkirtTriangles=0,VisibleSkirtTriangles=0;
+                    for(size_t Triangle=0;Triangle<Resident.Indices.size()/3;++Triangle)
+                    {
+                        const uint8 Faces=Triangle<Resident.TriangleSkirtFaceMasks.size()
+                            ?Resident.TriangleSkirtFaceMasks[Triangle]:0;
+                        if(Faces==0)++SurfaceTriangles;
+                        else
+                        {
+                            ++AllSkirtTriangles;
+                            if((Faces&Resident.ActiveSkirtFaceMask)!=0)++VisibleSkirtTriangles;
+                        }
+                    }
+                    std::printf("[CpuDrawOnlyTileAudit] key=(%d,%d,%d,lod%d) active_skirt_mask=0x%02x surface_triangles=%llu all_skirt_triangles=%llu visible_skirt_triangles=%llu draw_only_indices=%u\n",
+                        Tile.Coord.X,Tile.Coord.Y,Tile.Coord.Z,Tile.Level,
+                        static_cast<unsigned>(Resident.ActiveSkirtFaceMask),
+                        static_cast<unsigned long long>(SurfaceTriangles),
+                        static_cast<unsigned long long>(AllSkirtTriangles),
+                        static_cast<unsigned long long>(VisibleSkirtTriangles),
+                        static_cast<unsigned>(Indices.size()-TileIndexStart));
+                    bDrawOnlyCpuTriangleAuditReported=true;
+                }
                 if(bDrawOnlyTile&&IsDrawOnlyTile(Tile))
                 {
                     DrawOnlyIndexStart=static_cast<UINT>(TileIndexStart);
@@ -3736,6 +4284,14 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     static_cast<unsigned long long>(TotalIndices),StreamingBatches.size());
                 std::printf("[ColdTimeline] event=whole_clipmap_ready at_ms=%.3f waits_for=%zu_GPU_tiles_plus_CPU_fallback_and_uploads\n",
                     Elapsed*1000.0,GpuFarEligibleTiles.size());
+                if(GpuShaderWarmupThread.joinable()
+                    &&!bGpuShaderWarmupQueueFinished.load(std::memory_order_acquire))
+                {
+                    const bool bPriorityBoosted=SetThreadPriority(
+                        GpuShaderWarmupThread.native_handle(),THREAD_PRIORITY_HIGHEST)!=0;
+                    std::printf("[GpuShaderWarmupPriorityBoost] at_ms=%.3f reason=CPU_initial_clipmap_ready result=%s\n",
+                        Elapsed*1000.0,bPriorityBoosted?"highest":"failed");
+                }
                 std::printf("[PresentStartupSummary] nonblocking_startup_calls=%llu skipped_busy_calls=%llu gpu_warmup_held_calls=%llu policy=do_not_wait_until_initial_clipmap_ready\n",
                     static_cast<unsigned long long>(StartupNonBlockingPresentCalls),
                     static_cast<unsigned long long>(StartupNonBlockingPresentSkips),
@@ -3743,8 +4299,22 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 if(bGpuFarEnabled)
                 {
                     SampleGpuFarMemory();
+                    double CpuMeshWorkerSeconds=0.0;
+                    uint64 CpuMeshTiles=0;
+                    for(int32 Level=0;Level<=MaxLevels;++Level)
+                    {
+                        CpuMeshWorkerSeconds+=WorkerSecondsByLevel[Level];
+                        CpuMeshTiles+=GeneratedByLevel[Level];
+                    }
                     const double AverageGpuMs=GpuFarTilesSubmitted
                         ?GpuFarGpuDispatchMs/static_cast<double>(GpuFarTilesSubmitted):0.0;
+                    std::printf("[GpuFarModeTiming] mode=%s cpu_mesh_tiles=%llu cpu_mesh_worker_sum_s=%.6f cpu_input_export_worker_sum_s=%.6f gpu_submitted_tiles=%llu gpu_busy_timestamp_sum_ms=%.3f whole_view_wall_s=%.6f warmup_cpu_claimed=%llu warmup_gpu_takeover=%llu\n",
+                        bGpuFarAuditEnabled?"gpu_audit":"gpu_hybrid",
+                        static_cast<unsigned long long>(CpuMeshTiles),CpuMeshWorkerSeconds,
+                        GpuFarPreparationWorkerMs/1000.0,
+                        static_cast<unsigned long long>(GpuFarTilesSubmitted),GpuFarGpuDispatchMs,
+                        Elapsed,static_cast<unsigned long long>(GpuFarWarmupCpuClaimedTiles),
+                        static_cast<unsigned long long>(GpuFarWarmupGpuTakeoverTiles));
                     std::printf("[GpuFarTiming] submitted_tiles=%llu submission_batches=%llu max_tiles_per_batch=%u pooled_output_reserve_wall_ms=%.3f cpu_input_prepare_wall_ms=%.3f cpu_input_worker_sum_ms=%.3f cpu_upload_enqueue_sum_ms=%.3f gpu_dispatch_timestamp_total_ms=%.3f gpu_dispatch_mean_per_tile_ms=%.6f gpu_dispatch_max_tile_ms=%.6f gpu_submission_span_total_ms=%.3f gpu_submission_span_max_batch_ms=%.3f frame_budget_90hz_ms=11.111\n",
                         static_cast<unsigned long long>(GpuFarTilesSubmitted),
                         static_cast<unsigned long long>(GpuFarSubmissionBatches),GpuFarMaxTilesPerSubmissionBatch,
@@ -3767,8 +4337,25 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                         GpuFarLocalMemoryPeakBytes/(1024.0*1024.0),
                         GpuFarLocalMemoryBudgetBytes/(1024.0*1024.0));
                 }
+                else
+                {
+                    double CpuMeshWorkerSeconds=0.0;
+                    uint64 CpuMeshTiles=0;
+                    for(int32 Level=0;Level<=MaxLevels;++Level)
+                    {
+                        CpuMeshWorkerSeconds+=WorkerSecondsByLevel[Level];
+                        CpuMeshTiles+=GeneratedByLevel[Level];
+                    }
+                    std::printf("[GpuFarModeTiming] mode=cpu cpu_mesh_tiles=%llu cpu_mesh_worker_sum_s=%.6f whole_view_wall_s=%.6f\n",
+                        static_cast<unsigned long long>(CpuMeshTiles),CpuMeshWorkerSeconds,Elapsed);
+                }
                 ReportTileTimings();std::fflush(stdout);
-                if(bQuitAfterClipmap&&CapturePath.empty())PostMessageA(Window,WM_CLOSE,0,0);
+                if(bQuitAfterClipmap&&CapturePath.empty()
+                    &&(!bGpuFarEnabled||!bGpuFarShaderPrewarm||bGpuWarmupCompleted))
+                {
+                    bQuitAfterClipmapClosePosted=true;
+                    PostMessageA(Window,WM_CLOSE,0,0);
+                }
             }
             if(bInitialClipmapReported&&bGpuFarAuditEnabled&&GpuFarAudit&&!bGpuFarAuditReported)
             {
@@ -3910,7 +4497,28 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                         ||!Tile.VertexBuffer||!Tile.IndexBuffer||!Tile.IndirectArgs)continue;
                     if(!bInitialClipmapReported&&Pair.first.Level>=World->FullResClipLevels)continue;
                     if(!IsRenderedTileReady(Pair.first))continue;
-                    if(bAuditGpuSkirtDrawArgs&&Tile.VisibleSkirtFaceMask!=0&&!Tile.bSkirtDrawArgsAudited)
+                    const bool bDepthDisabledAudit=bGpuFarDepthDisabledAudit&&bDrawOnlyTile
+                        &&IsDrawOnlyTile(Pair.first);
+                    if(bDepthDisabledAudit)
+                    {
+                        Context->OMSetDepthStencilState(DepthDisabledAuditState,0);
+                        if(!bGpuFarDepthDisabledAuditReported)
+                        {
+                            bGpuFarDepthDisabledAuditReported=true;
+                            std::printf("[GpuFarDepthDisabledAudit] key=(%d,%d,%d,lod%d) depth_test=disabled path=gpu_tile_draw\n",
+                                Pair.first.Coord.X,Pair.first.Coord.Y,Pair.first.Coord.Z,Pair.first.Level);
+                        }
+                    }
+                    const bool bDrawReadbackCopy=(bGpuFarReadbackCopyDrawAudit||bGpuFarMeterCopyDrawAudit
+                        ||bGpuFarCpuMeshDrawAudit)
+                        &&bDrawOnlyTile
+                        &&IsDrawOnlyTile(Pair.first)&&Tile.ReadbackCopyVertexBuffer
+                        &&Tile.ReadbackCopyIndexBuffer;
+                    const bool bForceDirectDraw=bGpuFarDirectDrawAudit&&bDrawOnlyTile
+                        &&IsDrawOnlyTile(Pair.first);
+                    uint32 DirectDrawArgs[11]{};
+                    if(bAuditGpuSkirtDrawArgs&&!Tile.bSkirtDrawArgsAudited
+                        &&(Tile.VisibleSkirtFaceMask!=0||(bDrawOnlyTile&&IsDrawOnlyTile(Pair.first))))
                     {
                         if(!GpuSkirtArgsStaging)
                         {
@@ -3935,19 +4543,28 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                                 D3D11_MAP_READ,0,&ArgsMapped);
                             if(SUCCEEDED(MapResult))
                             {
-                                const auto* Args=static_cast<const uint32*>(ArgsMapped.pData);
-                                const uint32 ExpectedStartIndex=(Tile.IndirectArgsByteOffset
-                                    /(11u*sizeof(uint32)))*32768u;
+                            const auto* Args=static_cast<const uint32*>(ArgsMapped.pData);
+                            if(bForceDirectDraw)std::copy_n(Args,11,DirectDrawArgs);
+                            const uint32 LocalOutputSlot=Tile.IndirectArgsByteOffset
+                                    /(11u*sizeof(uint32));
+                                const uint32 ExpectedStartIndex=LocalOutputSlot*32768u;
+                                const uint32 ExpectedSurfaceStartIndex=LocalOutputSlot
+                                    *VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile;
                                 ++GpuSkirtDrawAuditTiles;
                                 GpuSkirtDrawAuditSumIndices+=Args[5];
                                 GpuSkirtDrawAuditMaxIndices=std::max<uint64>(GpuSkirtDrawAuditMaxIndices,Args[5]);
                                 if(Args[5]==0)++GpuSkirtDrawAuditZeroCounts;
                                 if(Args[10]!=0)++GpuSkirtDrawAuditOverflow;
                                 if(Args[7]!=ExpectedStartIndex)++GpuSkirtDrawAuditBadStartIndex;
-                                std::printf("[GpuSkirtDrawArgs] key=(%d,%d,%d,lod%d) face_mask=0x%02x output_slot=%u surface_indices=%u skirt_panel_indices=%u instances=%u connector_start_index=%u expected_start_index=%u overflow=%u\n",
+                                if(Args[0]%3u!=0||Args[0]>VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile
+                                    ||Args[1]!=1u||Args[2]!=ExpectedSurfaceStartIndex
+                                    ||Args[3]!=0u||Args[4]!=0u)
+                                    ++GpuSkirtDrawAuditBadSurfaceArgs;
+                                std::printf("[GpuSkirtDrawArgs] key=(%d,%d,%d,lod%d) face_mask=0x%02x output_slot=%u surface_indices=%u surface_instances=%u surface_start_index=%u expected_surface_start_index=%u surface_base_vertex=%u surface_start_instance=%u skirt_panel_indices=%u connector_start_index=%u expected_connector_start_index=%u overflow=%u\n",
                                     Pair.first.Coord.X,Pair.first.Coord.Y,Pair.first.Coord.Z,Pair.first.Level,
                                     static_cast<unsigned>(Tile.VisibleSkirtFaceMask),Tile.OutputSlot,
-                                    Args[0],Args[5],Args[6],Args[7],ExpectedStartIndex,Args[10]);
+                                    Args[0],Args[1],Args[2],ExpectedSurfaceStartIndex,Args[3],Args[4],
+                                    Args[5],Args[7],ExpectedStartIndex,Args[10]);
                                 Context->Unmap(GpuSkirtArgsStaging,0);
                             }
                             else std::fprintf(stderr,"[GpuSkirtDrawAuditError] map_staging_hr=0x%08lx\n",
@@ -3955,30 +4572,67 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                         }
                         Tile.bSkirtDrawArgsAudited=true;
                     }
-                    const UINT Stride=32,Offset=0;
-                    Context->IASetVertexBuffers(0,1,&Tile.VertexBuffer,&Stride,&Offset);
-                    Context->IASetIndexBuffer(Tile.IndexBuffer,DXGI_FORMAT_R32_UINT,0);
-                    Context->DrawIndexedInstancedIndirect(Tile.IndirectArgs,Tile.IndirectArgsByteOffset);
-                    if(Tile.ConnectorVertexBuffer&&Tile.ConnectorIndexBuffer)
+                    const UINT Stride=bDrawReadbackCopy?sizeof(FGpuVertex):32,Offset=0;
+                    if(bDrawReadbackCopy)
                     {
-                        Context->IASetVertexBuffers(0,1,&Tile.ConnectorVertexBuffer,&Stride,&Offset);
-                        Context->IASetIndexBuffer(Tile.ConnectorIndexBuffer,DXGI_FORMAT_R32_UINT,0);
-                        Context->DrawIndexedInstancedIndirect(Tile.IndirectArgs,
-                            Tile.IndirectArgsByteOffset+5u*sizeof(uint32));
+                        if(bGpuFarMeterCopyDrawAudit||bGpuFarCpuMeshDrawAudit)
+                        {
+                            C.Padding[0]=1.0f;
+                            Context->UpdateSubresource(Constants,0,nullptr,&C,0,0);
+                            Context->VSSetConstantBuffers(0,1,&Constants);
+                            Context->PSSetConstantBuffers(0,1,&Constants);
+                        }
+                        Context->IASetVertexBuffers(0,1,&Tile.ReadbackCopyVertexBuffer,&Stride,&Offset);
+                        Context->IASetIndexBuffer(Tile.ReadbackCopyIndexBuffer,DXGI_FORMAT_R32_UINT,0);
+                        Context->DrawIndexed(Tile.ReadbackCopyIndexCount,0,0);
                     }
+                    else
+                    {
+                        Context->IASetVertexBuffers(0,1,&Tile.VertexBuffer,&Stride,&Offset);
+                        Context->IASetIndexBuffer(Tile.IndexBuffer,DXGI_FORMAT_R32_UINT,0);
+                        if(bForceDirectDraw)
+                        {
+                            if(DirectDrawArgs[0]&&DirectDrawArgs[1])
+                                Context->DrawIndexed(DirectDrawArgs[0],DirectDrawArgs[2],
+                                    static_cast<INT>(DirectDrawArgs[3]));
+                            if(DirectDrawArgs[5]&&Tile.ConnectorVertexBuffer&&Tile.ConnectorIndexBuffer)
+                            {
+                                Context->IASetVertexBuffers(0,1,&Tile.ConnectorVertexBuffer,&Stride,&Offset);
+                                Context->IASetIndexBuffer(Tile.ConnectorIndexBuffer,DXGI_FORMAT_R32_UINT,0);
+                                Context->DrawIndexed(DirectDrawArgs[5],DirectDrawArgs[7],
+                                    static_cast<INT>(DirectDrawArgs[8]));
+                            }
+                            std::printf("[GpuDirectDrawAudit] key=(%d,%d,%d,lod%d) surface_indices=%u surface_start_index=%u surface_base_vertex=%u panel_plus_skirt_indices=%u connector_start_index=%u path=direct\n",
+                                Pair.first.Coord.X,Pair.first.Coord.Y,Pair.first.Coord.Z,Pair.first.Level,
+                                DirectDrawArgs[0],DirectDrawArgs[2],DirectDrawArgs[3],DirectDrawArgs[5],DirectDrawArgs[7]);
+                        }
+                        else
+                        {
+                            Context->DrawIndexedInstancedIndirect(Tile.IndirectArgs,Tile.IndirectArgsByteOffset);
+                            if(Tile.ConnectorVertexBuffer&&Tile.ConnectorIndexBuffer)
+                            {
+                                Context->IASetVertexBuffers(0,1,&Tile.ConnectorVertexBuffer,&Stride,&Offset);
+                                Context->IASetIndexBuffer(Tile.ConnectorIndexBuffer,DXGI_FORMAT_R32_UINT,0);
+                                Context->DrawIndexedInstancedIndirect(Tile.IndirectArgs,
+                                    Tile.IndirectArgsByteOffset+5u*sizeof(uint32));
+                            }
+                        }
+                    }
+                    if(bDepthDisabledAudit)Context->OMSetDepthStencilState(nullptr,0);
                 }
             }
             if(bAuditGpuSkirtDrawArgs&&bInitialClipmapReported&&!bGpuSkirtDrawAuditReported)
             {
-                std::printf("[GpuSkirtDrawArgsSummary] tiles=%llu zero_index_counts=%llu overflow=%llu bad_slot_start_indices=%llu total_panel_plus_skirt_indices=%llu max_panel_plus_skirt_indices=%llu status=%s\n",
+                std::printf("[GpuSkirtDrawArgsSummary] tiles=%llu zero_index_counts=%llu overflow=%llu bad_surface_args=%llu bad_connector_start_indices=%llu total_panel_plus_skirt_indices=%llu max_panel_plus_skirt_indices=%llu status=%s\n",
                     static_cast<unsigned long long>(GpuSkirtDrawAuditTiles),
                     static_cast<unsigned long long>(GpuSkirtDrawAuditZeroCounts),
                     static_cast<unsigned long long>(GpuSkirtDrawAuditOverflow),
+                    static_cast<unsigned long long>(GpuSkirtDrawAuditBadSurfaceArgs),
                     static_cast<unsigned long long>(GpuSkirtDrawAuditBadStartIndex),
                     static_cast<unsigned long long>(GpuSkirtDrawAuditSumIndices),
                     static_cast<unsigned long long>(GpuSkirtDrawAuditMaxIndices),
                     (GpuSkirtDrawAuditTiles>0&&GpuSkirtDrawAuditOverflow==0
-                        &&GpuSkirtDrawAuditBadStartIndex==0)?"PASS":"FAIL");
+                        &&GpuSkirtDrawAuditBadStartIndex==0&&GpuSkirtDrawAuditBadSurfaceArgs==0)?"PASS":"FAIL");
                 bGpuSkirtDrawAuditReported=true;
                 std::fflush(stdout);
             }
@@ -4009,6 +4663,26 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                         std::chrono::duration<double,std::milli>(GpuPumpEnd-CpuPumpEnd).count());
                 return;
             }
+            if(bGpuFarEnabled&&bGpuFarShaderPrewarm&&!bGpuWarmupCompleted
+                &&!bGpuShaderWarmupQueueFinished.load(std::memory_order_acquire))
+            {
+                if(!bGpuWarmupPresentSkipReported)
+                {
+                    bGpuWarmupPresentSkipReported=true;
+                    std::printf("[GpuWarmupPresentSkip] begin_ms=%.3f policy=skip_present_while_cold_shader_queue_is_pending\n",
+                        std::chrono::duration<double,std::milli>(
+                            std::chrono::steady_clock::now()-ColdProcessStart).count());
+                    std::fflush(stdout);
+                }
+                std::this_thread::sleep_for(std::chrono::milliseconds(1));
+                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 warmup_queue_pending=1\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();
             const bool bNonBlockingStartupPresent=bGpuFarEnabled&&!bInitialClipmapReported
                 &&!bForceBlockingStartupPresent;
@@ -4092,6 +4766,13 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     std::chrono::duration<double,std::milli>(BeforePresent-GpuPumpEnd).count(),
                     PresentCallMs,PresentSyncInterval,static_cast<unsigned long>(PresentResult));
             LastDrawEndAt=DrawEnd;
+            if(bQuitAfterClipmap&&CapturePath.empty()&&bInitialClipmapReported
+                &&(!bGpuFarEnabled||!bGpuFarShaderPrewarm||bGpuWarmupCompleted)
+                &&!bQuitAfterClipmapClosePosted)
+            {
+                bQuitAfterClipmapClosePosted=true;
+                PostMessageA(Window,WM_CLOSE,0,0);
+            }
         }
     };
 
@@ -5902,6 +6583,11 @@ int main(int Argc,char** Argv)
         else if(Arg=="--hold-cpu-frame-during-gpu-warmup")
             App.bHoldCpuFrameDuringGpuWarmup=true,App.bPresentCpuFrameBeforeGpu=true;
         else if(Arg=="--gpu-skirt-draw-audit")App.bAuditGpuSkirtDrawArgs=true;
+        else if(Arg=="--gpu-far-direct-draw-audit")App.bGpuFarDirectDrawAudit=true;
+        else if(Arg=="--gpu-far-readback-copy-draw-audit")App.bGpuFarReadbackCopyDrawAudit=true;
+        else if(Arg=="--gpu-far-meter-copy-draw-audit")App.bGpuFarMeterCopyDrawAudit=true;
+        else if(Arg=="--gpu-far-cpu-mesh-draw-audit")App.bGpuFarCpuMeshDrawAudit=true;
+        else if(Arg=="--gpu-far-depth-disabled-audit")App.bGpuFarDepthDisabledAudit=true;
         else if(Arg=="--gpu-skirts-disabled")App.bGpuSkirtEmitEnabled=false;
         else if(Arg=="--gpu-connectors-disabled")App.bGpuConnectorRepairEnabled=false;
         else if(Arg=="--profile-attribution")App.bProfileAttribution=true;
@@ -6197,6 +6883,11 @@ int main(int Argc,char** Argv)
             App.DataPath.c_str(),App.Position.X,App.Position.Y,App.Position.Z);
     }
     std::fflush(stdout);
+    if(App.bGpuFarEnabled&&(App.bGpuFarWarmupOnly||App.bGpuFarWarmupBeforeWorkers))
+    {
+        std::string GpuInitError;
+        if(!App.EnsureGpuFarPipeline(GpuInitError))return 3;
+    }
     if(App.bGpuFarWarmupOnly)
     {
         const auto WarmupOnlyStart=std::chrono::steady_clock::now();
@@ -6221,7 +6912,7 @@ int main(int Argc,char** Argv)
             std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-WarmupOnlyStart).count());
         return 0;
     }
-    App.StartGpuShaderWarmupAsync();
+    if(App.bGpuFarWarmupBeforeWorkers)App.StartGpuShaderWarmupAsync();
     if(App.bGpuFarWarmupBeforeWorkers)
     {
         if(App.GpuShaderWarmupThread.joinable())App.GpuShaderWarmupThread.join();
@@ -6231,6 +6922,8 @@ int main(int Argc,char** Argv)
         if(!VoxelForgeGpuFarViewer::CompleteFarShaderWarmup(App.Device,App.Context,
             App.GpuFarShaderPath,WarmupError))
         {std::fprintf(stderr,"[GpuShaderWarmupError] completion: %s\n",WarmupError.c_str());return 5;}
+        App.SetGpuWarmupContextProtection(false);
+        App.bGpuWarmupCompleted=true;
         std::printf("[GpuShaderWarmupBeforeWorkers] completed_before_generation_workers=true\n");
     }
     App.StartGeneration();
diff --git a/Standalone/Source/StreamingWorld.h b/Standalone/Source/StreamingWorld.h
index 723f8bd..b3ce256 100644
--- a/Standalone/Source/StreamingWorld.h
+++ b/Standalone/Source/StreamingWorld.h
@@ -478,7 +478,7 @@ public:
         for(std::thread& Worker:Workers)if(Worker.joinable())Worker.join();
         Workers.clear();
         std::lock_guard<std::mutex> Lock(Mutex);
-        Ready.clear();Outstanding.clear();Desired.clear();
+        Ready.clear();Outstanding.clear();Desired.clear();Running.clear();
     }
 
     FStandaloneStreamingUpdate UpdateDesired(
@@ -512,6 +512,75 @@ public:
         return Out;
     }
 
+    FStandaloneStreamingUpdate RestrictDesiredTileFilterKeepingStarted(
+        const std::vector<FVoxelTileKey>& BaseFilter,
+        const std::vector<FVoxelTileKey>& CandidateKeys,
+        const std::vector<FVoxelTileKey>& ResidentKeys,
+        std::vector<FVoxelTileKey>& RetainedCandidateKeys)
+    {
+        std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> Candidates;
+        Candidates.reserve(CandidateKeys.size()*2+1);
+        for(const FVoxelTileKey& Key:CandidateKeys)Candidates.insert(Key);
+
+        FStandaloneStreamingUpdate Stats;
+        std::lock_guard<std::mutex> Lock(Mutex);
+        std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> Retained;
+        Retained.reserve(CandidateKeys.size()*2+1);
+        for(const FVoxelTileKey& Key:Running)
+            if(Candidates.contains(Key))Retained.insert(Key);
+        for(const FStandaloneCpuTile& Tile:Ready)
+            if(Candidates.contains(Tile.Key))Retained.insert(Tile.Key);
+        for(const FVoxelTileKey& Key:ResidentKeys)
+            if(Candidates.contains(Key))Retained.insert(Key);
+
+        TileFilter.clear();
+        TileFilter.reserve((BaseFilter.size()+Retained.size())*2+1);
+        for(const FVoxelTileKey& Key:BaseFilter)TileFilter.insert(Key);
+        for(const FVoxelTileKey& Key:Retained)TileFilter.insert(Key);
+        bUseTileFilter=true;
+
+        FStandaloneDesiredTileSet NewDesired=CurrentDesired;
+        NewDesired.Tiles.erase(std::remove_if(NewDesired.Tiles.begin(),NewDesired.Tiles.end(),
+            [&](const FVoxelTileKey& Key){return !TileFilter.contains(Key);}),NewDesired.Tiles.end());
+        Stats.CenterChunk=NewDesired.CenterChunk;
+        Stats.DesiredTiles=static_cast<int32>(NewDesired.Tiles.size());
+        std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> ResidentSet;
+        ResidentSet.reserve(ResidentKeys.size()*2+1);
+        for(const FVoxelTileKey& Key:ResidentKeys)ResidentSet.insert(Key);
+        for(const FVoxelTileKey& Key:NewDesired.Tiles)
+        {
+            if(Key.Level<World->FullResClipLevels)++Stats.FullResTiles;
+            if(ResidentSet.contains(Key))++Stats.RetainedResidentTiles;
+        }
+
+        Desired.clear();Desired.reserve(NewDesired.Tiles.size()*2+1);
+        for(const FVoxelTileKey& Key:NewDesired.Tiles)Desired.insert(Key);
+        FullResTileCount=Stats.FullResTiles;
+        std::unordered_map<uint64,uint64> CancelledByCrossing;
+        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;
+        });
+        Queue.erase(QueueEnd,Queue.end());
+        std::make_heap(Queue.begin(),Queue.end(),FStandaloneTileJobHeapOrder{});
+        Stats.CancelledByCrossing.reserve(CancelledByCrossing.size());
+        for(const auto& Pair:CancelledByCrossing)Stats.CancelledByCrossing.push_back(Pair);
+        CurrentDesired=std::move(NewDesired);
+        RetainedCandidateKeys.assign(Retained.begin(),Retained.end());
+        std::sort(RetainedCandidateKeys.begin(),RetainedCandidateKeys.end(),
+            [](const FVoxelTileKey& A,const FVoxelTileKey& B)
+            {
+                if(A.Level!=B.Level)return A.Level<B.Level;
+                if(A.Coord.Z!=B.Coord.Z)return A.Coord.Z<B.Coord.Z;
+                if(A.Coord.Y!=B.Coord.Y)return A.Coord.Y<B.Coord.Y;
+                return A.Coord.X<B.Coord.X;
+            });
+        return Stats;
+    }
+
     bool IsDesired(const FVoxelTileKey& Key)const
     {std::lock_guard<std::mutex> Lock(Mutex);return Desired.contains(Key);}
     int32 DesiredCount()const
@@ -739,6 +808,7 @@ private:
                 if(bStop)return;
                 std::pop_heap(Queue.begin(),Queue.end(),FStandaloneTileJobHeapOrder{});
                 Job=Queue.back();Queue.pop_back();Job.StartedAt=FClock::now();
+                Running.insert(Job.Key);
             }
             if(bColdTimelineTrace)TraceTileEvent("started",Job.Key,Job.StartedAt,Job.QueuedAt,Job.StartedAt);
             FStandaloneCpuTile Tile=Generate(Job);
@@ -747,6 +817,7 @@ private:
             {
                 std::lock_guard<std::mutex> Lock(Mutex);
                 if(bStop)return;
+                Running.erase(Tile.Key);
                 Ready.push_back(std::move(Tile));
             }
         }
@@ -788,6 +859,7 @@ private:
     std::vector<std::thread> Workers;
     std::vector<FStandaloneTileJob> Queue;
     std::vector<FStandaloneCpuTile> Ready;
+    std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> Running;
     std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> Desired;
     std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> Outstanding;
     std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> TileFilter;