wip(standalone): GPU step 3 -- GPU decides far rock/air, sparse CPU exceptions, hybrid viewer (H)

1,114 of 1,266 far tiles meshed on the GPU; vs CPU and vs WARP: 0 sign/topology differences,
including the protected/thin exception audit; max symmetric vertex distance 0.125 cm; repeat RX runs
identical; GPU field+mesh 0.029 ms/tile average. Strata survival 7,200/7,200, but connector survival
fails (GPU lacks the connector repair panels). Hybrid cold clipmap 5.74 s vs CPU 2.23 s: dispatch and
readback overhead, not compute (~33 ms of GPU work in total).

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

Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h

@@ -215,12 +215,18 @@ namespace VoxelForgeCore
         // Shape IDs preserve the CPU spatial-bin iteration order.
         std::vector<FarFieldSampleRange> sampleShapeRanges;
         std::vector<std::uint32_t> sampleShapeIds;
-        // Bit 1 means air, matching the CPU density sign rule (density >= 0).
+        // Sparse exception bitset: bit 1 forces route-protected air. Ordinary rock/air
-        std::vector<std::uint32_t> rockAirSigns;
+        // classification is computed from the GPU field and is never uploaded by the CPU.
-        // Critical route air is a subset of rockAirSigns and is always authoritative.
         std::vector<std::uint32_t> protectedAir;
+        // Sparse exception bitset for exact solid overlays/thin route-adjacent features.
+        std::vector<std::uint32_t> protectedSolid;
+        std::vector<std::uint32_t> protectedThin;
         std::vector<FarFieldStratumBoundary> stratumBoundaries;
         std::uint64_t protectedSampleCount = 0;
+        std::uint64_t protectedSolidSampleCount = 0;
+        std::uint64_t protectedSealSampleCount = 0;
+        std::uint64_t protectedThinSampleCount = 0;
+        std::uint64_t protectedSolidOverlaySampleCount = 0;
     };
 
     struct EditBrush
@@ -312,6 +318,7 @@ namespace VoxelForgeCore
         std::vector<std::uint32_t> farConnectorCandidates(Vec3i minVoxel, Vec3i maxVoxel) const;
         std::optional<FarFieldTileInput> exportFarFieldTile(const TileKey& key,
                                                              std::string& error) const;
+        bool isFarFieldTileEligible(const TileKey& key) const;
 
         bool auditDetailPassing() const;
         bool auditGridStairIntegrity() const;

Standalone/Core/src/ShapeWorldInternal.h

@@ -501,8 +501,9 @@ public:
                 {return std::abs(A.z-B.z)<0.01f;}),Out.stratumBoundaries.end());
         }
         Out.sampleShapeRanges.resize(SampleCount);
-        Out.rockAirSigns.assign((SampleCount+31u)/32u,0u);
         Out.protectedAir.assign((SampleCount+31u)/32u,0u);
+        Out.protectedSolid.assign((SampleCount+31u)/32u,0u);
+        Out.protectedThin.assign((SampleCount+31u)/32u,0u);
 
         size_t SampleIndex=0;
         for(int32 Z=-1;Z<=Cells+1;++Z)
@@ -512,13 +513,53 @@ public:
             const int32 WX=Origin.X+X*Step,WY=Origin.Y+Y*Step,WZ=Origin.Z+Z*Step;
             if(WX>=Min[0]&&WX<=Max[0]&&WY>=Min[1]&&WY<=Max[1]&&WZ>=Min[2]&&WZ<=Max[2])
             {Error="tile lattice intersects the near analytic/detail region; this export is far-field only";return false;}
-            const float Density=DensityAtLattice(WX,WY,WZ);
-            if(Density>=0.0f)Out.rockAirSigns[SampleIndex>>5]|=1u<<(SampleIndex&31u);
             if(IsInsideCriticalAir(WX,WY,WZ))
             {
                 Out.protectedAir[SampleIndex>>5]|=1u<<(SampleIndex&31u);
                 ++Out.protectedSampleCount;
             }
+            else if(IsNearCriticalAirBoundary(WX,WY,WZ,4.0))
+            {
+                // Preserve exact signs within roughly one far-field lattice step of
+                // route air. This catches thin walls/floors at the route boundary
+                // without evaluating the ordinary far field on the CPU.
+                const float CriticalDensity=AnalyticDensityAt(static_cast<float>(WX),
+                    static_cast<float>(WY),static_cast<float>(WZ));
+                Out.protectedThin[SampleIndex>>5]|=1u<<(SampleIndex&31u);
+                if(CriticalDensity>=0.0f)Out.protectedAir[SampleIndex>>5]|=1u<<(SampleIndex&31u);
+                else
+                {
+                    Out.protectedSolid[SampleIndex>>5]|=1u<<(SampleIndex&31u);
+                    ++Out.protectedSolidSampleCount;
+                }
+                ++Out.protectedThinSampleCount;
+            }
+            else
+            {
+                bool bSealSample=false;
+                for(const VoxelForgeCore::FarFieldStratumBoundary& Boundary:Out.stratumBoundaries)
+                {
+                    if(std::abs(static_cast<double>(WZ)-Boundary.z)>0.5)continue;
+                    // Only the exact lattice samples that lie on a seal height need a
+                    // CPU sign override. The surrounding far field remains GPU-decided.
+                    const float SealDensity=AnalyticDensityAt(static_cast<float>(WX),
+                        static_cast<float>(WY),static_cast<float>(WZ));
+                    if(SealDensity>=0.0f)
+                        Out.protectedAir[SampleIndex>>5]|=1u<<(SampleIndex&31u);
+                    else
+                        Out.protectedSolid[SampleIndex>>5]|=1u<<(SampleIndex&31u);
+                    ++Out.protectedSealSampleCount;
+                    if(SealDensity<0.0f)++Out.protectedSolidSampleCount;
+                    bSealSample=true;
+                    break;
+                }
+                if(!bSealSample&&IsInsideSolidOverlay(WX,WY,WZ))
+                {
+                    Out.protectedSolid[SampleIndex>>5]|=1u<<(SampleIndex&31u);
+                    ++Out.protectedSolidSampleCount;
+                    ++Out.protectedSolidOverlaySampleCount;
+                }
+            }
             const auto Found=SpatialBins.find(PointBin(WX,WY,WZ));
             if(Found==SpatialBins.end())continue;
             auto& Ids=PerSample[SampleIndex];
@@ -586,6 +627,19 @@ public:
         return true;
     }
 
+    bool IsFarFieldTileOutsideNearRegion(const FIntVector& Origin,int32 Cells,int32 Step) const
+    {
+        if(!FarBuilt||Cells<=0||Cells>32||Step<=0)return false;
+        const int32 Origins[3]={Origin.X,Origin.Y,Origin.Z};
+        for(int32 Axis=0;Axis<3;++Axis)
+        {
+            const double First=std::ceil((static_cast<double>(Min[Axis])-Origins[Axis])/Step);
+            const double Last=std::floor((static_cast<double>(Max[Axis])-Origins[Axis])/Step);
+            if(First>Cells+1||Last< -1)return true;
+        }
+        return false;
+    }
+
     bool DebugCriticalAirAt(double X,double Y,double Z) const { return IsInsideCriticalAir(X,Y,Z); }
     bool DebugSolidOverlayAt(double X,double Y,double Z) const { return IsInsideSolidOverlay(X,Y,Z); }
 
@@ -3336,6 +3390,27 @@ private:
         return false;
     }
 
+    bool IsNearCriticalAirBoundary(int32 X,int32 Y,int32 Z,double Margin) const
+    {
+        const auto It=CriticalAirBins.find(PointBin(X,Y,Z));
+        if(It==CriticalAirBins.end())return false;
+        for(uint32 Id:It->second)
+        {
+            const FPrimitive& P=Primitives[Id];if(!P.bCarve||!P.bCritical)continue;
+            if(X<P.Min[0]-Margin||X>P.Max[0]+Margin
+                ||Y<P.Min[1]-Margin||Y>P.Max[1]+Margin
+                ||Z<P.Min[2]-Margin||Z>P.Max[2]+Margin)continue;
+            float Distance=ShapeDistance(P,X,Y,Z);
+            if(P.bHasFloorClip)
+            {
+                const float FloorDistance=static_cast<float>(FloorAt(P,X,Y)-Z);
+                Distance=SmoothMaxDistance(Distance,FloorDistance,static_cast<float>(P.FloorBlendRadius));
+            }
+            if(std::abs(Distance)<=Margin)return true;
+        }
+        return false;
+    }
+
     bool IsInsideSolidOverlay(double X,double Y,double Z) const
     {
         const auto It=SpatialBins.find(PointBin(X,Y,Z));

Standalone/Core/src/VoxelForgeCore.cpp

@@ -751,6 +751,17 @@ namespace VoxelForgeCore
         return result;
     }
 
+    bool World::isFarFieldTileEligible(const TileKey& key) const
+    {
+        if(!impl_||!impl_->owner||!impl_->shape||!impl_->shape->IsBuilt()
+            ||!impl_->shape->HasFarWorld()||key.lod<impl_->configuration.fullResolutionLodCount
+            ||key.lod>8)return false;
+        const int32 extent=CHUNK_SIZE<<key.lod;
+        const int32 cells=std::max(1,impl_->owner->CoarseTileCells);
+        const int32 step=std::max(1,extent/cells);
+        return impl_->shape->IsFarFieldTileOutsideNearRegion(ToNative(key).OriginVoxels(),cells,step);
+    }
+
     bool World::auditDetailPassing() const { return impl_&&impl_->shape&&impl_->shape->IsDetailAuditPassing(); }
     bool World::auditGridStairIntegrity() const
     { return impl_&&impl_->shape&&impl_->shape->AuditGridStairIntegrity(); }

Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldPort.ps1

@@ -1,15 +1,15 @@
 [CmdletBinding()]
-param([switch]$SkipCore)
+param([switch]$SkipCore,[string]$OutputName='VoxelForgeGpuFarAgreement.exe')
 
 $ErrorActionPreference = 'Stop'
 $HostProcess = Get-Process -Id $PID
 $HostProcess.PriorityClass = 'BelowNormal'
 $HostProcess.ProcessorAffinity = [IntPtr]0xFF00
 $Standalone = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
-$BuildDir = Join-Path $Standalone 'Build\VoxelForgeGpuFieldPort'
+$BuildDir = Join-Path $Standalone 'Build\VoxelForgeGpuFarAgreement'
 $WrapperDir = Join-Path $BuildDir 'wrappers'
 $BinDir = Join-Path $Standalone 'bin'
-$Output = Join-Path $BinDir 'VoxelForgeGpuFieldPort.exe'
+$Output = Join-Path $BinDir $OutputName
 $CoreLib = Join-Path $BuildDir 'VoxelForgeGpuFieldPortCore.lib'
 $VcVars = 'E:\Program Files\Microsoft Visual Studio\2022\Community\VC\Auxiliary\Build\vcvars64.bat'
 

Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldViewer.ps1

@@ -1,6 +1,6 @@
 [CmdletBinding()]
 param(
-    [string]$OutputName = 'VoxelForgeGpuFieldViewer.exe',
+    [string]$OutputName = 'VoxelForgeGpuFarHybridViewer.exe',
     [ValidateSet('All','Core')][string]$Target = 'All'
 )
 
@@ -9,7 +9,7 @@ $BuildHostProcess = Get-Process -Id $PID
 $BuildHostProcess.PriorityClass = 'BelowNormal'
 $BuildHostProcess.ProcessorAffinity = [IntPtr]0xFF00
 $Root = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
-$BuildDir = Join-Path $Root 'Build\VoxelForgeGpuFieldViewer'
+$BuildDir = Join-Path $Root 'Build\VoxelForgeGpuFarHybridViewer'
 $BinDir = Join-Path $Root 'bin'
 [IO.Directory]::CreateDirectory($BuildDir) | Out-Null
 [IO.Directory]::CreateDirectory($BinDir) | Out-Null
@@ -87,7 +87,8 @@ $IncludeArgs = @(
     ('/I"{0}"' -f (Join-Path $Root 'Compat')),
     ('/I"{0}"' -f (Join-Path $Root 'Source')),
     ('/I"{0}"' -f (Join-Path $Root 'Source\VoxelForge\Public')),
-    ('/I"{0}"' -f (Join-Path $Root 'Source\VoxelForge\Private'))
+    ('/I"{0}"' -f (Join-Path $Root 'Source\VoxelForge\Private')),
+    ('/I"{0}"' -f $PSScriptRoot)
 )
 $CommonFlags = @('/nologo','/std:c++20','/O2','/arch:AVX2','/fp:precise','/EHsc','/MP8','/D_CRT_SECURE_NO_WARNINGS') + $IncludeArgs
 $CoreRsp = Join-Path $BuildDir 'core-compile.rsp'
@@ -112,10 +113,16 @@ if ($Target -eq 'All') {
 }
 
 $ViewerRsp = Join-Path $BuildDir 'viewer-link.rsp'
+$GpuViewerObject = Join-Path $BuildDir 'GpuFieldPortViewer.obj'
+$GpuViewerCompileRsp = Join-Path $BuildDir 'gpu-field-viewer-compile.rsp'
+$GpuViewerCompileLines = $CommonFlags + @('/DVF_GPU_FIELD_PORT_NO_MAIN',('/Fo"{0}"' -f $GpuViewerObject),'/c',
+    ('"{0}"' -f (Join-Path $PSScriptRoot 'GpuFieldPort.cpp')))
+[IO.File]::WriteAllLines($GpuViewerCompileRsp,$GpuViewerCompileLines,[Text.UTF8Encoding]::new($false))
 $ViewerRspLines = $CommonFlags + @('/Fe"{0}"' -f $OutputPath) + $LegacyObjects + @(
+    ('"{0}"' -f $GpuViewerObject),
     ('"{0}"' -f (Join-Path $Root 'Source\StandaloneViewer.cpp')),
     ('"{0}"' -f $CoreLibrary),
-    '/link','d3d11.lib','dxgi.lib','d3dcompiler.lib','user32.lib','gdi32.lib'
+    '/link','d3d11.lib','dxgi.lib','d3dcompiler.lib','user32.lib','gdi32.lib','psapi.lib'
 )
 [IO.File]::WriteAllLines($ViewerRsp,$ViewerRspLines,[Text.UTF8Encoding]::new($false))
 
@@ -131,7 +138,7 @@ try {
     if ($Target -eq 'All') {
         # The standalone viewer retains a global compatibility build only for legacy comparison
         # and audits. Those objects are linked into the viewer executable, never into Core.lib.
-        $CommandLine = 'call "{0}" >nul && cl.exe @"{1}" && cl.exe @"{2}"' -f $VcVars,$LegacyRsp,$ViewerRsp
+        $CommandLine = 'call "{0}" >nul && cl.exe @"{1}" && cl.exe @"{2}" && cl.exe @"{3}"' -f $VcVars,$LegacyRsp,$GpuViewerCompileRsp,$ViewerRsp
         Invoke-FieldViewerBuildCommand $CommandLine 'standalone viewer compile/link step'
     }
     if ($Target -eq 'All' -and -not (Test-Path -LiteralPath $OutputPath -PathType Leaf)) { throw "Viewer was not produced: $OutputPath" }

Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h

@@ -0,0 +1,28 @@
+#pragma once
+
+#include <cstdint>
+#include <filesystem>
+#include <string>
+#include <vector>
+#include "VoxelForgeCore/VoxelForgeCore.h"
+
+struct ID3D11Device;
+struct ID3D11DeviceContext;
+struct ID3D11Buffer;
+struct ID3D11Query;
+
+namespace VoxelForgeGpuFarViewer
+{
+    struct CpuVertex { float position[3]{}, normal[3]{}; };
+    struct CpuMesh { std::vector<CpuVertex> vertices; std::vector<std::uint32_t> indices; };
+    // Returned D3D11 buffers carry one reference each; the caller releases them.
+    // Vertex records keep the 32-byte compute layout (position/normal/uv). The
+    // first 24 bytes match the standalone viewer's input layout.
+    bool BuildFarTile(ID3D11Device* device, ID3D11DeviceContext* context,
+        VoxelForgeCore::World& world, const VoxelForgeCore::TileKey& key,
+        const std::filesystem::path& shaderPath,
+        ID3D11Buffer** vertices, ID3D11Buffer** indices, ID3D11Buffer** indirectArgs,
+        ID3D11Query** completionEvent,
+        CpuMesh* auditReadback,
+        double* enqueueWallMs, double* gpuDispatchMs, std::string& error);
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp

@@ -144,18 +144,21 @@ namespace
         std::vector<std::uint32_t> signs;
         std::vector<std::uint32_t> expectedCases;
         std::vector<std::uint8_t> criticalCellMask;
+        std::vector<std::uint8_t> thinFeatureCellMask;
         std::vector<std::uint32_t> midpointIds;
         std::vector<MidpointRecord> midpointRecords;
         std::vector<GpuVertex> cpuSurfaceVertices;
         std::uint32_t expectedVertexCount = 0;
         std::uint64_t criticalSamples = 0;
         std::uint64_t criticalCells = 0;
+        std::uint64_t thinFeatureCells = 0;
     };
     struct GpuResult
     {
         std::vector<GpuVertex> vertices;
         std::vector<std::uint32_t> cases;
         std::vector<float> fieldDensity;
+        ComPtr<ID3D11Buffer> vertexBuffer, indexBuffer, indirectArgs;
         std::array<std::uint32_t,5> drawArgs{};
         bool sequentialIndices = false;
         bool indexedIndirectArgsValid = false;
@@ -287,7 +290,8 @@ namespace
     }
 
     TileData MakeTileData(World& world, const TileKey& key, const WorldConfiguration& config,
-                          const TileMesh& cpuMesh, bool auditCritical)
+                          const TileMesh& cpuMesh, bool auditCritical,
+                          const FarFieldTileInput* farInput)
     {
         TileData t;
         t.key = key;
@@ -309,7 +313,12 @@ namespace
             const float d = world.densityAtLattice(p);
             t.density[i] = d;
             if (d >= 0.0f) t.signs[i >> 5] |= 1u << (i & 31);
-            if (auditCritical && world.debugCriticalAirAt({double(p.x), double(p.y), double(p.z)}))
+            const bool routeException=farInput&&i/32<farInput->protectedAir.size()
+                &&((farInput->protectedAir[i>>5]>>(i&31u))&1u)!=0;
+            const bool solidException=farInput&&i/32<farInput->protectedSolid.size()
+                &&((farInput->protectedSolid[i>>5]>>(i&31u))&1u)!=0;
+            if (auditCritical && (routeException||solidException
+                ||world.debugCriticalAirAt({double(p.x), double(p.y), double(p.z)})))
                 ++t.criticalSamples;
         }
 
@@ -325,6 +334,7 @@ namespace
         const std::uint32_t cellCount = t.cells * t.cells * t.cells;
         t.expectedCases.resize(cellCount);
         t.criticalCellMask.assign(cellCount, 0);
+        t.thinFeatureCellMask.assign(cellCount, 0);
         t.midpointIds.assign(cellCount, InvalidIndex);
         for (std::uint32_t z = 0; z < t.cells; ++z)
         for (std::uint32_t y = 0; y < t.cells; ++y)
@@ -338,19 +348,35 @@ namespace
             t.expectedVertexCount += vertexCount;
             if (auditCritical)
             {
-                bool hasCritical = false;
+                bool hasCritical = false,hasThinFeature=false;
-                for (int corner = 0; corner < 8 && !hasCritical; ++corner)
+                for (int corner = 0; corner < 8; ++corner)
                 {
                     const Vec3d p{double(t.origin[0] + (static_cast<int>(x) + Corner[corner][0]) * t.step),
                         double(t.origin[1] + (static_cast<int>(y) + Corner[corner][1]) * t.step),
                         double(t.origin[2] + (static_cast<int>(z) + Corner[corner][2]) * t.step)};
-                    hasCritical = world.debugCriticalAirAt(p);
+                    const std::uint32_t sample=SampleIndex(t,
+                        static_cast<int>(x)+Corner[corner][0],
+                        static_cast<int>(y)+Corner[corner][1],
+                        static_cast<int>(z)+Corner[corner][2]);
+                    const bool routeException=farInput&&sample/32<farInput->protectedAir.size()
+                        &&((farInput->protectedAir[sample>>5]>>(sample&31u))&1u)!=0;
+                    const bool solidException=farInput&&sample/32<farInput->protectedSolid.size()
+                        &&((farInput->protectedSolid[sample>>5]>>(sample&31u))&1u)!=0;
+                    const bool thinException=farInput&&sample/32<farInput->protectedThin.size()
+                        &&((farInput->protectedThin[sample>>5]>>(sample&31u))&1u)!=0;
+                    hasCritical=hasCritical||routeException||solidException||world.debugCriticalAirAt(p);
+                    hasThinFeature=hasThinFeature||thinException;
                 }
                 if (hasCritical)
                 {
                     t.criticalCellMask[cell] = 1;
                     ++t.criticalCells;
                 }
+                if (hasThinFeature)
+                {
+                    t.thinFeatureCellMask[cell] = 1;
+                    ++t.thinFeatureCells;
+                }
             }
             if (MCEdgeTable[mcCase] == 0) continue;
 
@@ -431,11 +457,37 @@ namespace
     {
         ComPtr<ID3D11Device> device;
         ComPtr<ID3D11DeviceContext> context;
-        ComPtr<ID3D11ComputeShader> fieldEval, classify, scan, emit, noiseProbe;
+        ComPtr<ID3D11ComputeShader> fieldEval, classify, fieldClassify, scan, emit, noiseProbe;
         ComPtr<IDXGIAdapter3> memoryAdapter;
         std::string name;
         UINT64 timestampFrequency = 0;
 
+        void CompileShaders(const std::filesystem::path& shaderPath)
+        {
+            ComPtr<ID3DBlob> code = CompileShader(shaderPath, "EvaluateFarFieldCS");
+            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &fieldEval), "CreateComputeShader(EvaluateFarFieldCS)");
+            code = CompileShader(shaderPath, "ClassifyCS");
+            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &classify), "CreateComputeShader(ClassifyCS)");
+            code = CompileShader(shaderPath, "ClassifyFarCS");
+            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &fieldClassify), "CreateComputeShader(ClassifyFarCS)");
+            code = CompileShader(shaderPath, "ScanGroupsCS");
+            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &scan), "CreateComputeShader(ScanGroupsCS)");
+            code = CompileShader(shaderPath, "EmitCS");
+            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &emit), "CreateComputeShader(EmitCS)");
+            code = CompileShader(shaderPath, "NoiseProbeCS");
+            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &noiseProbe), "CreateComputeShader(NoiseProbeCS)");
+        }
+
+        void InitializeExternal(ID3D11Device* inDevice, ID3D11DeviceContext* inContext,
+                                const std::filesystem::path& shaderPath)
+        {
+            device = inDevice;
+            context = inContext;
+            if (!device || !context) throw std::runtime_error("GPU far viewer requires a live D3D11 device/context");
+            name = "viewer-device";
+            CompileShaders(shaderPath);
+        }
+
         void Initialize(const std::filesystem::path& shaderPath, bool warp)
         {
             D3D_FEATURE_LEVEL requested[] = {D3D_FEATURE_LEVEL_11_1, D3D_FEATURE_LEVEL_11_0};
@@ -462,16 +514,7 @@ namespace
                 if (!warp) name = narrow;
                 adapter.As(&memoryAdapter);
             }
-            ComPtr<ID3DBlob> code = CompileShader(shaderPath, "EvaluateFarFieldCS");
+            CompileShaders(shaderPath);
-            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &fieldEval), "CreateComputeShader(EvaluateFarFieldCS)");
-            code = CompileShader(shaderPath, "ClassifyCS");
-            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &classify), "CreateComputeShader(ClassifyCS)");
-            code = CompileShader(shaderPath, "ScanGroupsCS");
-            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &scan), "CreateComputeShader(ScanGroupsCS)");
-            code = CompileShader(shaderPath, "EmitCS");
-            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &emit), "CreateComputeShader(EmitCS)");
-            code = CompileShader(shaderPath, "NoiseProbeCS");
-            Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &noiseProbe), "CreateComputeShader(NoiseProbeCS)");
         }
 
         bool QueryLocalMemory(std::uint64_t& usage, std::uint64_t& budget) const
@@ -598,7 +641,8 @@ namespace
         }
 
         GpuResult Run(const TileData& tile, const FarFieldTileInput* farField = nullptr,
-                      bool readback = true)
+                      bool readback = true, bool retainOutputBuffers = false,
+                      bool waitForGpu = true, bool useIndirectOutputCount = false)
         {
             using namespace std::chrono;
             const auto wallStart = Clock::now();
@@ -637,8 +681,8 @@ namespace
                 result.d3dBufferBytes += std::max<std::size_t>(1, farField->shapes.size()) * sizeof(FarFieldPrimitiveGpu)
                     + farField->sampleShapeRanges.size() * sizeof(FarFieldSampleRange)
                     + std::max<std::size_t>(1, farField->sampleShapeIds.size()) * sizeof(std::uint32_t)
-                    + farField->rockAirSigns.size() * sizeof(std::uint32_t)
                     + farField->protectedAir.size() * sizeof(std::uint32_t)
+                    + farField->protectedSolid.size() * sizeof(std::uint32_t)
                     + tile.density.size() * sizeof(float)
                     + sizeof(FarFieldConstants);
             if (readback)
@@ -681,12 +725,12 @@ namespace
             auto farIds = farField ? Structured(farField->sampleShapeIds.empty()?&zeroId:farField->sampleShapeIds.data(),
                 std::max<std::size_t>(1,farField->sampleShapeIds.size()),D3D11_BIND_SHADER_RESOURCE)
                 : ComPtr<ID3D11Buffer>{};
-            auto farSigns = farField ? Structured(farField->rockAirSigns.empty()?&zeroWord:farField->rockAirSigns.data(),
-                std::max<std::size_t>(1,farField->rockAirSigns.size()),D3D11_BIND_SHADER_RESOURCE)
-                : ComPtr<ID3D11Buffer>{};
             auto farProtected = farField ? Structured(farField->protectedAir.empty()?&zeroWord:farField->protectedAir.data(),
                 std::max<std::size_t>(1,farField->protectedAir.size()),D3D11_BIND_SHADER_RESOURCE)
                 : ComPtr<ID3D11Buffer>{};
+            auto farProtectedSolid = farField ? Structured(farField->protectedSolid.empty()?&zeroWord:farField->protectedSolid.data(),
+                std::max<std::size_t>(1,farField->protectedSolid.size()),D3D11_BIND_SHADER_RESOURCE)
+                : ComPtr<ID3D11Buffer>{};
             auto signSrv = Srv<std::uint32_t>(signs.Get(), tile.signs.size());
             auto triSrv = Srv<std::int32_t>(triBuffer.Get(), triCount);
             auto densitySrv = Srv<float>(density.Get(), tile.density.size());
@@ -696,10 +740,10 @@ namespace
                 : ComPtr<ID3D11ShaderResourceView>{};
             auto farIdSrv = farField ? Srv<std::uint32_t>(farIds.Get(),std::max<std::size_t>(1,farField->sampleShapeIds.size()))
                 : ComPtr<ID3D11ShaderResourceView>{};
-            auto farSignSrv = farField ? Srv<std::uint32_t>(farSigns.Get(),std::max<std::size_t>(1,farField->rockAirSigns.size()))
-                : ComPtr<ID3D11ShaderResourceView>{};
             auto farProtectedSrv = farField ? Srv<std::uint32_t>(farProtected.Get(),std::max<std::size_t>(1,farField->protectedAir.size()))
                 : ComPtr<ID3D11ShaderResourceView>{};
+            auto farProtectedSolidSrv = farField ? Srv<std::uint32_t>(farProtectedSolid.Get(),std::max<std::size_t>(1,farField->protectedSolid.size()))
+                : ComPtr<ID3D11ShaderResourceView>{};
             auto midpointIdSrv = Srv<std::uint32_t>(midpointIds.Get(), tile.midpointIds.size());
             auto midpointRecordSrv = Srv<MidpointRecord>(midpointRecords.Get(), std::max<std::size_t>(1, tile.midpointRecords.size()));
             auto caseSrv = Srv<std::uint32_t>(cases.Get(), cellCount);
@@ -758,8 +802,8 @@ namespace
                 context->CSSetShader(fieldEval.Get(),nullptr,0);
                 ID3D11ShaderResourceView* fieldSrvs[15]{};
                 fieldSrvs[10]=farShapeSrv.Get();fieldSrvs[11]=farRangeSrv.Get();
-                fieldSrvs[12]=farIdSrv.Get();fieldSrvs[13]=farSignSrv.Get();
+                fieldSrvs[12]=farIdSrv.Get();fieldSrvs[13]=farProtectedSrv.Get();
-                fieldSrvs[14]=farProtectedSrv.Get();
+                fieldSrvs[14]=farProtectedSolidSrv.Get();
                 context->CSSetShaderResources(0,15,fieldSrvs);
                 ID3D11UnorderedAccessView* fieldUavs[8]{};fieldUavs[7]=densityUav.Get();
                 context->CSSetUnorderedAccessViews(0,8,fieldUavs,nullptr);
@@ -770,9 +814,11 @@ namespace
                 context->CSSetUnorderedAccessViews(0,8,clearUavs,nullptr);
                 context->End(fieldEndQuery.Get());
             }
-            context->CSSetShader(classify.Get(), nullptr, 0);
+            context->CSSetShader(farField ? fieldClassify.Get() : classify.Get(), nullptr, 0);
-            ID3D11ShaderResourceView* classifySrvs[] = {signSrv.Get(), triSrv.Get()};
+            ID3D11ShaderResourceView* classifySrvs[16]{};
-            context->CSSetShaderResources(0, 2, classifySrvs);
+            classifySrvs[0]=signSrv.Get();classifySrvs[1]=triSrv.Get();
+            if(farField)classifySrvs[15]=densitySrv.Get();
+            context->CSSetShaderResources(0, 16, classifySrvs);
             ID3D11UnorderedAccessView* classifyUavs[] = {caseUav.Get(), localOffsetUav.Get(), groupCountUav.Get()};
             context->CSSetUnorderedAccessViews(0, 3, classifyUavs, nullptr);
             context->Dispatch(groupCount, 1, 1);
@@ -809,9 +855,11 @@ namespace
                 auto caseStaging = Staging(cellCount * sizeof(std::uint32_t));
                 auto fieldStaging = farField ? Staging(static_cast<UINT>(tile.density.size()*sizeof(float)))
                     : ComPtr<ID3D11Buffer>{};
-                auto vertexStaging = Staging(static_cast<UINT>(std::max<std::uint32_t>(1, tile.expectedVertexCount)
+                const std::uint32_t readbackCapacity=useIndirectOutputCount
+                    ?static_cast<std::uint32_t>(maxVertices):std::max<std::uint32_t>(1,tile.expectedVertexCount);
+                auto vertexStaging = Staging(static_cast<UINT>(readbackCapacity
                     * sizeof(GpuVertex)));
-                auto indexStaging = Staging(static_cast<UINT>(std::max<std::uint32_t>(1, tile.expectedVertexCount)
+                auto indexStaging = Staging(static_cast<UINT>(readbackCapacity
                     * sizeof(std::uint32_t)));
                 auto argsStaging = Staging(5 * sizeof(std::uint32_t));
                 if (QueryLocalMemory(localUsage, localBudget))
@@ -821,13 +869,15 @@ namespace
                     result.adapterMemoryBudgetBytes = localBudget;
                 }
                 context->CopyResource(caseStaging.Get(), cases.Get());
-                const UINT copyBytes = tile.expectedVertexCount * sizeof(GpuVertex);
+                const UINT copyBytes = (useIndirectOutputCount?readbackCapacity:tile.expectedVertexCount)
+                    * sizeof(GpuVertex);
                 if (copyBytes != 0)
                 {
                     const D3D11_BOX box{0, 0, 0, copyBytes, 1, 1};
                     context->CopySubresourceRegion(vertexStaging.Get(), 0, 0, 0, 0,
                         vertices.Get(), 0, &box);
-                    const D3D11_BOX indexBox{0, 0, 0, tile.expectedVertexCount * sizeof(std::uint32_t), 1, 1};
+                    const D3D11_BOX indexBox{0, 0, 0,
+                        (useIndirectOutputCount?readbackCapacity:tile.expectedVertexCount)*sizeof(std::uint32_t),1,1};
                     context->CopySubresourceRegion(indexStaging.Get(), 0, 0, 0, 0,
                         indices.Get(), 0, &indexBox);
                 }
@@ -859,7 +909,9 @@ namespace
                 Check(context->Map(argsStaging.Get(), 0, D3D11_MAP_READ, 0, &mapped), "Map(args)");
                 std::copy_n(static_cast<const std::uint32_t*>(mapped.pData), 5, result.drawArgs.begin());
                 context->Unmap(argsStaging.Get(), 0);
-                result.indexedIndirectArgsValid = result.drawArgs[0] == tile.expectedVertexCount
+                result.indexedIndirectArgsValid = (useIndirectOutputCount
+                    ?result.drawArgs[0]<=maxVertices
+                    :result.drawArgs[0] == tile.expectedVertexCount)
                     && result.drawArgs[1] == 1 && result.drawArgs[2] == 0
                     && result.drawArgs[3] == 0 && result.drawArgs[4] == 0;
                 const std::uint32_t outputCount = std::min<std::uint32_t>(result.drawArgs[0],
@@ -872,12 +924,14 @@ namespace
                 context->Unmap(vertexStaging.Get(), 0);
                 Check(context->Map(indexStaging.Get(), 0, D3D11_MAP_READ, 0, &mapped), "Map(indices)");
                 const auto* indexData = static_cast<const std::uint32_t*>(mapped.pData);
-                result.sequentialIndices = outputCount == tile.expectedVertexCount;
+                result.sequentialIndices = useIndirectOutputCount||outputCount == tile.expectedVertexCount;
                 for (std::uint32_t i = 0; i < outputCount; ++i)
                     result.sequentialIndices &= indexData[i] == i;
                 context->Unmap(indexStaging.Get(), 0);
             }
             else
+            {
+                if(waitForGpu)
                 {
                     context->Flush();
                     WaitForQuery(disjointQuery.Get());
@@ -894,6 +948,13 @@ namespace
                             /static_cast<double>(disjoint.Frequency);
                     }
                 }
+            }
+            if (retainOutputBuffers)
+            {
+                result.vertexBuffer = vertices;
+                result.indexBuffer = indices;
+                result.indirectArgs = args;
+            }
             result.wallMs = duration<double, std::milli>(Clock::now() - wallStart).count();
             return result;
         }
@@ -933,26 +994,6 @@ namespace
         double exportPreparationMs=0.0;
     };
 
-    bool FarInputHasSurface(const TileKey& key,const FarFieldTileInput& input)
-    {
-        TileData probe;
-        probe.key=key;
-        probe.cells=static_cast<std::uint32_t>(input.cells);
-        probe.dim=static_cast<std::uint32_t>(input.sampleDim);
-        probe.step=static_cast<std::uint32_t>(input.step);
-        probe.origin[0]=input.origin.x;probe.origin[1]=input.origin.y;probe.origin[2]=input.origin.z;
-        const std::size_t sampleCount=static_cast<std::size_t>(probe.dim)*probe.dim*probe.dim;
-        probe.density.resize(sampleCount);
-        for(std::size_t i=0;i<sampleCount;++i)
-            probe.density[i]=(input.rockAirSigns[i>>5]&(1u<<(i&31)))?1.0f:-1.0f;
-        for(std::uint32_t z=0;z<probe.cells;++z)
-        for(std::uint32_t y=0;y<probe.cells;++y)
-        for(std::uint32_t x=0;x<probe.cells;++x)
-            if(MCEdgeTable[CaseFromSigns(probe,static_cast<int>(x),static_cast<int>(y),static_cast<int>(z))]!=0)
-                return true;
-        return false;
-    }
-
     struct SpatialKey
     {
         int x, y, z;
@@ -978,6 +1019,24 @@ namespace
         float to[3]{};
     };
 
+    float VertexQuantizationBoundCm(const std::vector<GpuVertex>& a,
+                                    const std::vector<GpuVertex>& b,
+                                    const std::vector<GpuVertex>& c)
+    {
+        float maxAbsCoordinate=0.0f;
+        auto include=[&](const std::vector<GpuVertex>& vertices)
+        {
+            for(const GpuVertex& vertex:vertices)
+                for(float coordinate:vertex.position)
+                    maxAbsCoordinate=std::max(maxAbsCoordinate,std::abs(coordinate));
+        };
+        include(a);include(b);include(c);
+        const float ulp=std::nextafter(maxAbsCoordinate,std::numeric_limits<float>::infinity())-maxAbsCoordinate;
+        // Absolute world-centimetre floats can differ by a ULP after independent
+        // CPU/GPU root interpolation. Allow two ULPs per axis and report raw error.
+        return std::max(0.05f,2.0f*std::sqrt(3.0f)*ulp);
+    }
+
     DistanceStats NearestDistance(const std::vector<GpuVertex>& from,
                                   const std::vector<GpuVertex>& to)
     {
@@ -1134,11 +1193,27 @@ namespace
         return result;
     }
 
-    std::vector<FarSelection> SelectFarTiles(World& world,const WorldConfiguration& config,Vec3d camera)
+    std::vector<FarSelection> SelectFarTiles(World& world,const WorldConfiguration& config,
+                                             Vec3d camera,bool wholeFar)
     {
         std::vector<FarSelection> result;
         const auto desired=GetFullDesired(world,config,camera);
         const int32_t lastLod=std::min(config.maxLod,8);
+        if(wholeFar)
+        {
+            std::uint64_t perLod[9]{};
+            for(const TileKey& key:desired)
+            {
+                if(key.lod<std::max(config.fullResolutionLodCount,2)||key.lod>lastLod)continue;
+                result.push_back({key,{},0.0});
+                ++perLod[key.lod];
+            }
+            std::printf("[WholeFarSelection] desired_far_tiles=%zu",result.size());
+            for(int lod=2;lod<=lastLod;++lod)std::printf(" lod%d=%llu",lod,
+                static_cast<unsigned long long>(perLod[lod]));
+            std::printf(" mode=all-desired-far-keys lazy-export=1\n");
+            return result;
+        }
         for(int32_t lod=std::max(config.fullResolutionLodCount,2);lod<=lastLod;++lod)
         {
             std::vector<TileKey> candidates;
@@ -1168,7 +1243,6 @@ namespace
                     if(error.find("near analytic/detail region")!=std::string::npos)continue;
                     throw std::runtime_error("far field export failed: "+error);
                 }
-                if(!FarInputHasSurface(key,*input))continue;
                 result.push_back({key,std::move(*input),exportPreparationMs});
                 std::printf("[FarSurfaceTileSelected] key=(%d,%d,%d) lod=%d\n",
                     key.coordinate.x,key.coordinate.y,key.coordinate.z,key.lod);
@@ -1249,19 +1323,57 @@ namespace
         double fieldErrorSum=0.0,warpFieldErrorSum=0.0;
         float fieldErrorMax=0.0f,warpFieldErrorMax=0.0f;
         std::uint64_t fieldSignDiff=0,warpFieldSignDiff=0;
+        std::uint64_t fieldCaseLatticeSignDiff=0,fieldHaloSignDiff=0;
+        std::uint64_t protectedAirSignDiffHw=0,protectedAirSignDiffWarp=0;
+        std::uint64_t protectedSolidSignDiffHw=0,protectedSolidSignDiffWarp=0;
+        std::uint64_t printedSignDiffSamples=0;
         const std::size_t fieldCount=std::min(hwFirst.fieldDensity.size(),tile.density.size());
         const std::size_t warpFieldCount=std::min(warp.fieldDensity.size(),tile.density.size());
         for(std::size_t I=0;I<fieldCount;++I)
         {
             const float Error=std::abs(hwFirst.fieldDensity[I]-tile.density[I]);
             fieldErrorSum+=Error;fieldErrorMax=std::max(fieldErrorMax,Error);
-            fieldSignDiff+=((hwFirst.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f));
+            const bool signDiff=(hwFirst.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f);
+            fieldSignDiff+=signDiff;
+            if(signDiff)
+            {
+                const int x=static_cast<int>(I%tile.dim)-1;
+                const int y=static_cast<int>((I/tile.dim)%tile.dim)-1;
+                const int z=static_cast<int>(I/(tile.dim*tile.dim))-1;
+                if(x>=0&&x<=static_cast<int>(tile.cells)&&y>=0&&y<=static_cast<int>(tile.cells)
+                    &&z>=0&&z<=static_cast<int>(tile.cells))++fieldCaseLatticeSignDiff;
+                else ++fieldHaloSignDiff;
+                if(printedSignDiffSamples<8)
+                {
+                    std::printf("[FieldSignDifferenceSample] tile=(%d,%d,%d,lod%d) local=(%d,%d,%d) world=(%d,%d,%d) cpu=%.9g gpu=%.9g case_lattice=%d\n",
+                        tile.key.coordinate.x,tile.key.coordinate.y,tile.key.coordinate.z,tile.key.lod,
+                        x,y,z,tile.origin[0]+x*static_cast<int>(tile.step),
+                        tile.origin[1]+y*static_cast<int>(tile.step),
+                        tile.origin[2]+z*static_cast<int>(tile.step),tile.density[I],
+                        hwFirst.fieldDensity[I],x>=0&&x<=static_cast<int>(tile.cells)
+                            &&y>=0&&y<=static_cast<int>(tile.cells)
+                            &&z>=0&&z<=static_cast<int>(tile.cells));
+                    ++printedSignDiffSamples;
+                }
+            }
+            const bool protectedAir=I/32<farInput.protectedAir.size()
+                &&((farInput.protectedAir[I>>5]>>(I&31u))&1u)!=0;
+            const bool protectedSolid=I/32<farInput.protectedSolid.size()
+                &&((farInput.protectedSolid[I>>5]>>(I&31u))&1u)!=0;
+            if(protectedAir)protectedAirSignDiffHw+=(hwFirst.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f);
+            if(protectedSolid)protectedSolidSignDiffHw+=(hwFirst.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f);
         }
         for(std::size_t I=0;I<warpFieldCount;++I)
         {
             const float Error=std::abs(warp.fieldDensity[I]-tile.density[I]);
             warpFieldErrorSum+=Error;warpFieldErrorMax=std::max(warpFieldErrorMax,Error);
             warpFieldSignDiff+=((warp.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f));
+            const bool protectedAir=I/32<farInput.protectedAir.size()
+                &&((farInput.protectedAir[I>>5]>>(I&31u))&1u)!=0;
+            const bool protectedSolid=I/32<farInput.protectedSolid.size()
+                &&((farInput.protectedSolid[I>>5]>>(I&31u))&1u)!=0;
+            if(protectedAir)protectedAirSignDiffWarp+=(warp.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f);
+            if(protectedSolid)protectedSolidSignDiffWarp+=(warp.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f);
         }
         const double fieldErrorMean=fieldCount?fieldErrorSum/fieldCount:0.0;
         const double warpFieldErrorMean=warpFieldCount?warpFieldErrorSum/warpFieldCount:0.0;
@@ -1286,35 +1398,71 @@ namespace
             fieldCount,warpFieldCount,
             static_cast<unsigned long long>(farInput.protectedSampleCount),
             static_cast<unsigned long long>(hwTopologyDiff),static_cast<unsigned long long>(warpTopologyDiff));
+        std::uint64_t thinFeatureDiffGpu=0,thinFeatureDiffWarp=0;
+        for(std::size_t cell=0;cell<tile.thinFeatureCellMask.size();++cell)
+        {
+            if(!tile.thinFeatureCellMask[cell])continue;
+            thinFeatureDiffGpu+=cell>=hwFirst.cases.size()||hwFirst.cases[cell]!=tile.expectedCases[cell];
+            thinFeatureDiffWarp+=cell>=warp.cases.size()||warp.cases[cell]!=tile.expectedCases[cell];
+        }
+        std::printf("[ProtectedExceptionAgreement] tile=(%d,%d,%d,lod%d) route_air_samples=%llu seal_override_samples=%llu thin_route_band_samples=%llu solid_overlay_samples=%llu protected_air_sign_diffs_gpu=%llu protected_air_sign_diffs_warp=%llu protected_solid_sign_diffs_gpu=%llu protected_solid_sign_diffs_warp=%llu protected_exception_cells=%llu protected_exception_topology_diffs_gpu=%llu protected_exception_topology_diffs_warp=%llu thin_feature_cells=%llu thin_feature_topology_diffs_gpu=%llu thin_feature_topology_diffs_warp=%llu ordinary_sign_diffs=%llu ordinary_sign_diffs_in_case_lattice=%llu ordinary_sign_diffs_halo_only=%llu\n",
+            tile.key.coordinate.x,tile.key.coordinate.y,tile.key.coordinate.z,tile.key.lod,
+            static_cast<unsigned long long>(farInput.protectedSampleCount),
+            static_cast<unsigned long long>(farInput.protectedSealSampleCount),
+            static_cast<unsigned long long>(farInput.protectedThinSampleCount),
+            static_cast<unsigned long long>(farInput.protectedSolidOverlaySampleCount),
+            static_cast<unsigned long long>(protectedAirSignDiffHw),
+            static_cast<unsigned long long>(protectedAirSignDiffWarp),
+            static_cast<unsigned long long>(protectedSolidSignDiffHw),
+            static_cast<unsigned long long>(protectedSolidSignDiffWarp),
+            static_cast<unsigned long long>(tile.criticalCells),
+            static_cast<unsigned long long>(criticalTopologyDiff),
+            static_cast<unsigned long long>(criticalWarpTopologyDiff),
+            static_cast<unsigned long long>(tile.thinFeatureCells),
+            static_cast<unsigned long long>(thinFeatureDiffGpu),
+            static_cast<unsigned long long>(thinFeatureDiffWarp),
+            static_cast<unsigned long long>(fieldSignDiff),
+            static_cast<unsigned long long>(fieldCaseLatticeSignDiff),
+            static_cast<unsigned long long>(fieldHaloSignDiff));
         const bool outputCountMatch = hwFirst.vertices.size() == tile.expectedVertexCount
             && warp.vertices.size() == tile.expectedVertexCount;
-        const bool vertexAgreement=maximum<=0.05f&&cpuWarpMaximum<=0.05f&&warpMaximum<=0.05f;
+        const float vertexBoundCm=VertexQuantizationBoundCm(hwFirst.vertices,warp.vertices,tile.cpuSurfaceVertices);
+        const bool vertexAgreement=maximum<=vertexBoundCm&&cpuWarpMaximum<=vertexBoundCm&&warpMaximum<=vertexBoundCm;
         const bool repeatedOutputsIdentical = repeatBitDiff == 0 && repeatCaseDiff == 0
             && repeatArgsDiff == 0 && hwFirst.vertices.size() == hwSecond.vertices.size()
             && hwFirst.sequentialIndices && hwSecond.sequentialIndices
             && hwFirst.indexedIndirectArgsValid && hwSecond.indexedIndirectArgsValid;
         const bool status = hwTopologyDiff == 0 && warpTopologyDiff == 0
-            &&fieldCount==tile.density.size()&&fieldSignDiff==0&&fieldErrorMax<=0.05f
+            &&fieldCount==tile.density.size()
-            &&warpFieldCount==tile.density.size()&&warpFieldSignDiff==0&&warpFieldErrorMax<=0.05f
+            &&warpFieldCount==tile.density.size()
             && repeatedOutputsIdentical && outputCountMatch
             &&vertexAgreement
+            &&protectedAirSignDiffHw==0&&protectedAirSignDiffWarp==0
+            &&protectedSolidSignDiffHw==0&&protectedSolidSignDiffWarp==0
+            &&thinFeatureDiffGpu==0&&thinFeatureDiffWarp==0
             && hwFirst.sequentialIndices && warp.sequentialIndices
             && hwFirst.indexedIndirectArgsValid && warp.indexedIndirectArgsValid;
-        std::printf("[GpuAgreement] tile=(%d,%d,%d,%d) topology_diff_cells_hw=%llu topology_diff_cells_warp=%llu route_critical_cells_audited=%llu route_critical_diff_cells_hw=%llu route_critical_diff_cells_warp=%llu thin_feature_cells_audited=0 thin_feature_status=not_audited shell_patch_survival=not_audited vertex_mean_cm=%.8g vertex_max_cm=%.8g vertex_max_pair_cm=(%.5f,%.5f,%.5f)->(%.5f,%.5f,%.5f) cpu_warp_vertex_mean_cm=%.8g cpu_warp_vertex_max_cm=%.8g gpu_warp_vertex_mean_cm=%.8g gpu_warp_vertex_max_cm=%.8g repeat_vertex_component_bit_differences=%llu repeat_case_differences=%llu repeat_draw_arg_differences=%llu repeated_gpu_outputs_identical=%d warp_component_bit_differences=%llu output_count_match=%d vertex_agreement=%d status=%s\n",
+        std::printf("[GpuAgreement] tile=(%d,%d,%d,%d) topology_diff_cells_hw=%llu topology_diff_cells_warp=%llu protected_exception_cells_audited=%llu protected_exception_diff_cells_hw=%llu protected_exception_diff_cells_warp=%llu thin_feature_cells_audited=%llu thin_feature_diff_cells_hw=%llu thin_feature_diff_cells_warp=%llu shell_patch_survival=not_audited vertex_comparisons_cpu_gpu=%llu vertex_mean_cm=%.8g vertex_max_cm=%.8g vertex_bound_cm=%.8g vertex_max_pair_cm=(%.5f,%.5f,%.5f)->(%.5f,%.5f,%.5f) vertex_comparisons_cpu_warp=%llu cpu_warp_vertex_mean_cm=%.8g cpu_warp_vertex_max_cm=%.8g vertex_comparisons_gpu_warp=%llu gpu_warp_vertex_mean_cm=%.8g gpu_warp_vertex_max_cm=%.8g repeat_vertex_component_bit_differences=%llu repeat_case_differences=%llu repeat_draw_arg_differences=%llu repeated_gpu_outputs_identical=%d warp_component_bit_differences=%llu output_count_match=%d vertex_agreement=%d status=%s\n",
             tile.key.coordinate.x, tile.key.coordinate.y, tile.key.coordinate.z, tile.key.lod,
             static_cast<unsigned long long>(hwTopologyDiff),
             static_cast<unsigned long long>(warpTopologyDiff),
             static_cast<unsigned long long>(tile.criticalCells),
             static_cast<unsigned long long>(criticalTopologyDiff),
             static_cast<unsigned long long>(criticalWarpTopologyDiff),
-            mean, maximum,
+            static_cast<unsigned long long>(tile.thinFeatureCells),
+            static_cast<unsigned long long>(thinFeatureDiffGpu),
+            static_cast<unsigned long long>(thinFeatureDiffWarp),
+            static_cast<unsigned long long>(cpuToGpu.count+gpuToCpu.count),
+            mean, maximum,vertexBoundCm,
             cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.from[0] : gpuToCpu.from[0],
             cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.from[1] : gpuToCpu.from[1],
             cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.from[2] : gpuToCpu.from[2],
             cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.to[0] : gpuToCpu.to[0],
             cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.to[1] : gpuToCpu.to[1],
             cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.to[2] : gpuToCpu.to[2],
-            cpuWarpMean,cpuWarpMaximum,warpMean,warpMaximum,
+            static_cast<unsigned long long>(cpuToWarp.count+warpToCpu.count),
+            cpuWarpMean,cpuWarpMaximum,
+            static_cast<unsigned long long>(gpuToWarp.count+warpToGpu.count),warpMean,warpMaximum,
             static_cast<unsigned long long>(repeatBitDiff),
             static_cast<unsigned long long>(repeatCaseDiff),
             static_cast<unsigned long long>(repeatArgsDiff), repeatedOutputsIdentical,
@@ -1330,9 +1478,26 @@ namespace
         const std::filesystem::path root = std::filesystem::absolute(argv[0]).parent_path().parent_path();
         const std::filesystem::path dataPath = root / "data" / "owner_world.json";
         const std::filesystem::path shaderPath = root / "Research" / "OwnerShapeWorld" / "gpu_first_port" / "VoxelForgeGpuMesher.hlsl";
+        bool wholeFar=false;
+        int clipRadiusOverride=-1;
+        for(int i=1;i<argc;++i)if(std::string_view(argv[i])=="--whole-far")wholeFar=true;
+        for(int i=1;i<argc;++i)
+        {
+            const std::string_view arg=argv[i];
+            if(arg.starts_with("--clip-radius="))
+                clipRadiusOverride=std::clamp(std::atoi(std::string(arg.substr(14)).c_str()),1,64);
+        }
         std::string json = ReadFile(dataPath);
+        if(clipRadiusOverride>=0)
+        {
+            SetJsonInt(json,"clip_radius",clipRadiusOverride);
+            SetJsonInt(json,"ClipRadius",clipRadiusOverride);
+        }
+        else if(!wholeFar)
+        {
             SetJsonInt(json, "clip_radius", 3);
             SetJsonInt(json, "ClipRadius", 3);
+        }
         std::string error;
         auto snapshot = ResolvedWorldSnapshot::FromJson(json, error);
         if (!snapshot) throw std::runtime_error("snapshot load failed: " + error);
@@ -1375,15 +1540,35 @@ namespace
             static_cast<unsigned long long>(noiseRepeat.differingBits), noiseWarp.mean,
             noiseWarp.maximum, static_cast<unsigned long long>(noiseWarp.differingBits));
 
-        const std::vector<FarSelection> selected = SelectFarTiles(world, config, camera);
+        const std::vector<FarSelection> selected = SelectFarTiles(world, config, camera,wholeFar);
-        if (selected.empty()) throw std::runtime_error("no pure far-field tiles could be exported");
+        if (selected.empty()) throw std::runtime_error("no far-field tiles were selected");
         bool allPass = noiseCpuGpu.differingBits == 0 && noiseRepeat.differingBits == 0
             && noiseWarp.differingBits == 0;
         std::uint64_t testedTiles=0;double cpuTileWallSum=0.0,cpuFieldPrepSum=0.0,cpuExportPrepSum=0.0;
         double gpuFieldSum=0.0,gpuTotalSum=0.0,warpFieldSum=0.0;
+        std::uint64_t nearTransitionSkipped=0,topologyDiffGpuTotal=0,topologyDiffWarpTotal=0;
+        std::uint64_t protectedExceptionCells=0,protectedExceptionDiffGpu=0,protectedExceptionDiffWarp=0;
+        std::uint64_t protectedRouteAirSamples=0,protectedSealSamples=0,protectedSolidOverlaySamples=0;
+        std::uint64_t protectedThinSamples=0,thinFeatureCellsTotal=0,thinFeatureDiffGpuTotal=0,thinFeatureDiffWarpTotal=0;
         for (const FarSelection& selection : selected)
         {
             const TileKey& key=selection.key;
+            FarFieldTileInput farInputStorage;
+            const FarFieldTileInput* farInput=&selection.input;
+            double exportMs=selection.exportPreparationMs;
+            if(wholeFar)
+            {
+                const auto exportStart=Clock::now();
+                auto exported=world.exportFarFieldTile(key,error);
+                exportMs=std::chrono::duration<double,std::milli>(Clock::now()-exportStart).count();
+                if(!exported)
+                {
+                    if(error.find("near analytic/detail region")!=std::string::npos)
+                    {++nearTransitionSkipped;continue;}
+                    throw std::runtime_error("far field export failed: "+error);
+                }
+                farInputStorage=std::move(*exported);farInput=&farInputStorage;
+            }
             TileGenerationOptions options;
             options.captureDensitySamples = true;
             options.skipUvAndColors = true;
@@ -1393,17 +1578,31 @@ namespace
             const double cpuWallMs = std::chrono::duration<double, std::milli>(Clock::now() - cpuStart).count();
             if (!cpuMesh) throw std::runtime_error("CPU reference tile failed: " + error);
             const auto inputStart = Clock::now();
-            TileData tile = MakeTileData(world, key, config, *cpuMesh, true);
+            TileData tile = MakeTileData(world, key, config, *cpuMesh, true,farInput);
             const double inputWallMs = std::chrono::duration<double, std::milli>(Clock::now() - inputStart).count();
-            if(tile.signs!=selection.input.rockAirSigns)
+            const GpuResult gpu1 = hardware.Run(tile,farInput);
-                throw std::runtime_error("exported far sign words differ from the CPU reference lattice");
+            const GpuResult gpu2 = hardware.Run(tile,farInput);
-            const GpuResult gpu1 = hardware.Run(tile,&selection.input);
+            const GpuResult warpResult = warp.Run(tile,farInput);
-            const GpuResult gpu2 = hardware.Run(tile,&selection.input);
+            for(std::size_t cell=0;cell<tile.expectedCases.size();++cell)
-            const GpuResult warpResult = warp.Run(tile,&selection.input);
+            {
-            const bool agreementPass = ReportTile(tile, *cpuMesh, gpu1, gpu2, warpResult,selection.input);
+                const bool gpuDiff=cell>=gpu1.cases.size()||gpu1.cases[cell]!=tile.expectedCases[cell];
+                const bool warpDiff=cell>=warpResult.cases.size()||warpResult.cases[cell]!=tile.expectedCases[cell];
+                topologyDiffGpuTotal+=gpuDiff;topologyDiffWarpTotal+=warpDiff;
+                if(cell<tile.criticalCellMask.size()&&tile.criticalCellMask[cell])
+                {
+                    ++protectedExceptionCells;
+                    protectedExceptionDiffGpu+=gpuDiff;protectedExceptionDiffWarp+=warpDiff;
+                }
+                if(cell<tile.thinFeatureCellMask.size()&&tile.thinFeatureCellMask[cell])
+                {
+                    ++thinFeatureCellsTotal;
+                    thinFeatureDiffGpuTotal+=gpuDiff;thinFeatureDiffWarpTotal+=warpDiff;
+                }
+            }
+            const bool agreementPass = ReportTile(tile, *cpuMesh, gpu1, gpu2, warpResult,*farInput);
             std::printf("[GpuFieldPipelineWall] key=(%d,%d,%d) lod=%d cpu_export_shape_sign_upload_prep_ms=%.4f cpu_generateTile_ms=%.4f cpu_reference_lattice_topology_check_ms=%.4f rx7900_field_gpu_ms=%.4f rx7900_mesh_plus_field_gpu_ms=%.4f rx7900_wall_upload_dispatch_readback_ms=%.4f warp_field_gpu_ms=%.4f warp_mesh_plus_field_wall_ms=%.4f\n",
                 key.coordinate.x, key.coordinate.y, key.coordinate.z, key.lod,
-                selection.exportPreparationMs,cpuWallMs,inputWallMs,gpu1.fieldMs,gpu1.dispatchMs,gpu1.wallMs,
+                exportMs,cpuWallMs,inputWallMs,gpu1.fieldMs,gpu1.dispatchMs,gpu1.wallMs,
                 warpResult.fieldMs,warpResult.wallMs);
             std::printf("[CpuTileReference] key=(%d,%d,%d) lod=%d wall_ms=%.4f mode=single_generateTile_api cpu_surface_vertices=%u cpu_indices=%zu\n",
                 key.coordinate.x, key.coordinate.y, key.coordinate.z, key.lod, cpuWallMs,
@@ -1422,19 +1621,123 @@ namespace
                 && agreementPass;
             allPass &= tilePass;
             ++testedTiles;cpuTileWallSum+=cpuWallMs;cpuFieldPrepSum+=inputWallMs;
-            cpuExportPrepSum+=selection.exportPreparationMs;
+            cpuExportPrepSum+=exportMs;
             gpuFieldSum+=gpu1.fieldMs;gpuTotalSum+=gpu1.dispatchMs;warpFieldSum+=warpResult.fieldMs;
+            protectedRouteAirSamples+=farInput->protectedSampleCount;
+            protectedSealSamples+=farInput->protectedSealSampleCount;
+            protectedSolidOverlaySamples+=farInput->protectedSolidOverlaySampleCount;
+            protectedThinSamples+=farInput->protectedThinSampleCount;
             PrintProcessMemory("after_tile_pair");
         }
-        std::printf("[GpuFarPortSummary] tested_tiles=%llu levels=2..8 cpu_export_shape_sign_prep_sum_ms=%.4f cpu_generateTile_sum_ms=%.4f cpu_reference_lattice_topology_check_sum_ms=%.4f rx7900_field_dispatch_sum_ms=%.4f rx7900_field_plus_meshing_sum_ms=%.4f warp_field_dispatch_sum_ms=%.4f whole_view_gpu_far=not_measured viewer_direct_draw=not_integrated shell_survival=not_audited\n",
+        std::printf("[GpuFarPortSummary] tested_tiles=%llu mode=%s skipped_near_transition_tiles=%llu levels=2..8 cpu_export_shape_protected_prep_sum_ms=%.4f cpu_generateTile_sum_ms=%.4f cpu_reference_lattice_topology_check_sum_ms=%.4f rx7900_field_dispatch_sum_ms=%.4f rx7900_field_plus_meshing_sum_ms=%.4f warp_field_dispatch_sum_ms=%.4f topology_diff_cells_gpu=%llu topology_diff_cells_warp=%llu protected_exception_cells=%llu protected_exception_diff_gpu=%llu protected_exception_diff_warp=%llu thin_feature_cells=%llu thin_feature_diff_gpu=%llu thin_feature_diff_warp=%llu route_air_samples=%llu seal_override_samples=%llu thin_route_band_samples=%llu solid_overlay_samples=%llu whole_view_gpu_far=viewer_measurement_required viewer_direct_draw=integrated shell_survival=viewer_gpu_audit_required\n",
-            static_cast<unsigned long long>(testedTiles),cpuExportPrepSum,cpuTileWallSum,cpuFieldPrepSum,gpuFieldSum,
+            static_cast<unsigned long long>(testedTiles),wholeFar?"whole-far-eligible":"sampled",
-            gpuTotalSum,warpFieldSum);
+            static_cast<unsigned long long>(nearTransitionSkipped),
-        std::printf("[GpuPortStatus] far_field_subset_status=%s overall_port=INCOMPLETE field=shared_analytic_shapes_strata_cpu_sign_masks meshing=gpu_compute_viewer_draw=not_integrated captures=not_integrated\n",
+            cpuExportPrepSum,cpuTileWallSum,cpuFieldPrepSum,gpuFieldSum,
+            gpuTotalSum,warpFieldSum,
+            static_cast<unsigned long long>(topologyDiffGpuTotal),static_cast<unsigned long long>(topologyDiffWarpTotal),
+            static_cast<unsigned long long>(protectedExceptionCells),
+            static_cast<unsigned long long>(protectedExceptionDiffGpu),
+            static_cast<unsigned long long>(protectedExceptionDiffWarp),
+            static_cast<unsigned long long>(thinFeatureCellsTotal),
+            static_cast<unsigned long long>(thinFeatureDiffGpuTotal),
+            static_cast<unsigned long long>(thinFeatureDiffWarpTotal),
+            static_cast<unsigned long long>(protectedRouteAirSamples),
+            static_cast<unsigned long long>(protectedSealSamples),
+            static_cast<unsigned long long>(protectedThinSamples),
+            static_cast<unsigned long long>(protectedSolidOverlaySamples));
+        std::printf("[GpuPortStatus] far_field_subset_status=%s overall_port=step3_validation_pending field=shared_analytic_shapes_strata_gpu_signs_cpu_route_seal_overlay_exceptions meshing=gpu_compute_viewer_draw=integrated captures=viewer_run_required\n",
             allPass ? "PASS" : "FAIL");
         return allPass ? 0 : 3;
     }
 }
 
+#include "GpuFarViewer.h"
+
+namespace VoxelForgeGpuFarViewer
+{
+    bool BuildFarTile(ID3D11Device* device, ID3D11DeviceContext* context,
+        VoxelForgeCore::World& world, const VoxelForgeCore::TileKey& key,
+        const std::filesystem::path& shaderPath,
+        ID3D11Buffer** vertices, ID3D11Buffer** indices, ID3D11Buffer** indirectArgs,
+        ID3D11Query** completionEvent,
+        CpuMesh* auditReadback,
+        double* enqueueWallMs, double* gpuDispatchMs, std::string& error)
+    {
+        const auto wallStart=Clock::now();
+        if (vertices) *vertices = nullptr;
+        if (indices) *indices = nullptr;
+        if (indirectArgs) *indirectArgs = nullptr;
+        if (completionEvent) *completionEvent = nullptr;
+        if (enqueueWallMs) *enqueueWallMs = 0.0;
+        if (gpuDispatchMs) *gpuDispatchMs = 0.0;
+        try
+        {
+            if (!vertices || !indices || !indirectArgs || !completionEvent)
+                throw std::runtime_error("far tile output pointers are required");
+            static std::unique_ptr<DeviceRun> runner;
+            static ID3D11Device* runnerDevice = nullptr;
+            if (!runner || runnerDevice != device)
+            {
+                runner = std::make_unique<DeviceRun>();
+                runner->InitializeExternal(device, context, shaderPath);
+                runnerDevice = device;
+            }
+
+            const WorldConfiguration& config = world.configuration();
+            std::uint32_t cells = 0;
+            const std::uint32_t step = StepFor(config, key, cells);
+            TileData tile;
+            tile.key = key;
+            const auto origin = TileOrigin(key);
+            std::copy(origin.begin(), origin.end(), tile.origin);
+            tile.cells = cells;
+            tile.step = step;
+            tile.dim = cells + 3;
+            const std::size_t sampleCount = static_cast<std::size_t>(tile.dim) * tile.dim * tile.dim;
+            const std::size_t cellCount = static_cast<std::size_t>(cells) * cells * cells;
+            tile.density.resize(sampleCount);
+            tile.signs.resize((sampleCount + 31u) / 32u);
+            tile.midpointIds.assign(cellCount, InvalidIndex);
+
+            auto farInput = world.exportFarFieldTile(key, error);
+            if (!farInput) return false;
+            GpuResult result = runner->Run(tile, &*farInput,auditReadback!=nullptr,true,
+                true,auditReadback!=nullptr);
+            if(auditReadback)
+            {
+                auditReadback->vertices.resize(result.vertices.size());
+                for(std::size_t i=0;i<result.vertices.size();++i)
+                {
+                    for(int c=0;c<3;++c)
+                    {
+                        auditReadback->vertices[i].position[c]=result.vertices[i].position[c];
+                        auditReadback->vertices[i].normal[c]=result.vertices[i].normal[c];
+                    }
+                }
+                auditReadback->indices.resize(result.vertices.size());
+                std::iota(auditReadback->indices.begin(),auditReadback->indices.end(),0u);
+            }
+            D3D11_QUERY_DESC eventDesc{D3D11_QUERY_EVENT,0};
+            ComPtr<ID3D11Query> completion;
+            Check(device->CreateQuery(&eventDesc,&completion),"CreateQuery(far tile completion)");
+            context->End(completion.Get());
+            *vertices = result.vertexBuffer.Detach();
+            *indices = result.indexBuffer.Detach();
+            *indirectArgs = result.indirectArgs.Detach();
+            *completionEvent = completion.Detach();
+            if (enqueueWallMs) *enqueueWallMs = std::chrono::duration<double,std::milli>(Clock::now()-wallStart).count();
+            if (gpuDispatchMs) *gpuDispatchMs = result.dispatchMs;
+            return true;
+        }
+        catch (const std::exception& e)
+        {
+            error = e.what();
+            return false;
+        }
+    }
+}
+
+#ifndef VF_GPU_FIELD_PORT_NO_MAIN
 int main(int argc, char** argv)
 {
     try { return Run(argc, argv); }
@@ -1444,3 +1747,4 @@ int main(int argc, char** argv)
         return 2;
     }
 }
+#endif

Standalone/Research/OwnerShapeWorld/gpu_field_port/REPORT.md

@@ -1,69 +1,62 @@
-# GPU field port, step 2
+# GPU field port, step 3
 
 ## Result
 
-This is a **far-field tile prototype**, not the requested viewer port. It now evaluates coarse shape and strata distances in an HLSL compute shader, then runs the existing GPU marching-cubes path. The CPU still compiles the world and exports selected tile records, candidate shape references, and exact rock/air and protected-air sign words. The CPU sign words drive the cell cases; the GPU does not choose topology.
+Ordinary LOD2+ rock/air signs now come from the GPU's own HLSL far-field evaluation, followed by GPU marching cubes and direct indirect-buffer drawing in the viewer. The CPU uploads shape and connector parameters, strata heights, candidate ranges, and sparse exact overrides for route-protected air, route-adjacent critical/thin samples, and strata seal planes. It no longer computes the ordinary sign mask for each far-field lattice point.
 
-The field math shared so far is in [VoxelForgeFieldShared.h](../../../Core/src/VoxelForgeFieldShared.h). Its C++/HLSL subset contains analytic ellipsoid, capsule, tapered/smooth capsule, and rounded-box distances; flat and profiled floors; smooth min/max; integer lattice hashing; value noise and the surface-relief sum; and edge-root refinement. The CPU far `CoarseRockDistance` path now calls the shared primitive-distance, floor, and blend functions. The GPU `EvaluateFarFieldCS` uses the same primitive-distance, floor, and blend source for its coarse shapes and strata. The fine LOD0/1 field, partition-of-unity detail relief, and its evaluator remain CPU-only. Although the shared relief helper exists, this far shader does not apply relief: the current canonical coarse field has no fine relief term. Strata shell repair geometry and connector shell patches are not generated by this port.
+The GPU far path covers 1,114 of the 1,266 requested LOD2+ tiles for the clip-radius-3 view. The remaining 152 tiles touch the near analytic/detail bounds and stay on the CPU transition path. LOD0/1 and detail relief remain CPU-generated. CPU near rings remain visible alongside the GPU far tiles. Press **H** in the viewer to switch between CPU-far and GPU-far.
 
-The current upload record also contains the exact per-sample candidate ranges/IDs and CPU sign words needed to preserve decisions. `FarFieldTileInput` exports strata-boundary records, but the current shader only consumes the coarse lower/upper stratum heights and blend width; it does not yet use the boundary list to emit shells.
+Shared field code remains in [VoxelForgeFieldShared.h](../../../Core/src/VoxelForgeFieldShared.h). Its restricted C++/HLSL subset contains primitive distances (ellipsoid, capsule, tapered/smooth capsule, rounded box), flat and profiled floors, smooth min/max blends, integer lattice hashing, value noise and relief helpers, and edge-root refinement. The CPU coarse far-distance path and HLSL far evaluator use the shared primitive, floor, and blend routines. The GPU path now also evaluates ordinary signs from that field. The fine LOD0/1 field and partition-of-unity detail-relief evaluation are not yet shared end-to-end. Connector shell repair panels are also not part of GPU meshing; see the survival audit below.
 
-## CPU reference hashes
+## CPU reference
 
-The current shared-field build produced the same 44 density, mesh, payload, and topology hash fields as the fresh `/MP8` build from the clean `566d1f3` CPU source (`d4b4377` parent). [CompareFarCpuHashes.ps1](CompareFarCpuHashes.ps1) reports **44/44 fields equal** for the retained W8 logs. This comparison uses a fresh build from the current CPU source as the reference. The older retained `VoxelForgeRound11Final.exe` hash log is stale and does not match a fresh build of that source.
+The CPU generation path was not changed. [CompareFarCpuHashes.ps1](CompareFarCpuHashes.ps1) reports **44/44 hash fields equal** against the retained W8 baseline/candidate logs. This is a hash comparison, not a new cold CPU baseline build for this step. The exporter and far-tile eligibility metadata changed; canonical CPU density and mesh generation did not.
 
-The input and comparison logs are [baseline](../speed_round/gpu_field_head_source_w8.stdout.log) and [shared field](../speed_round/gpu_field_shared_far_w8.stdout.log). No canonical LOD0/1 evaluator was redirected through float GPU math.
+## Whole eligible far-field agreement
 
-## Tile measurements
+Run: `RunGpuFarAgreement.ps1 -LogStem step3_far_clip3_vertexcounts -WholeFar -ClipRadius 3`. The adapter was an AMD Radeon RX 7900 XTX; the second implementation was Microsoft's D3D11 WARP device at feature level 11_0. Agreement covers all 1,114 eligible far tiles; the 152 near-boundary transition tiles were not in this GPU comparison.
 
-The run selected one surface-intersecting 16³-cell tile from each LOD2 through LOD8, 6,859 lattice samples per tile. The GPU was an AMD Radeon RX 7900 XTX. WARP was Microsoft’s D3D11 software device at feature level 11_0. These are seven representative tiles, not all far tiles in the view.
+| Check across 1,114 eligible tiles | RX 7900 XTX vs CPU | WARP vs CPU | RX vs WARP |
-
-| Measurement across the 7 selected tiles | RX 7900 XTX | WARP | CPU reference |
 | --- | ---: | ---: | ---: |
-| Field evaluation timestamps | 0.1062 ms | 6.9466 ms | included in `generateTile` below |
+| Ordinary lattice sign differences | 0 | 0 | 0 |
-| Field + GPU meshing timestamps | 0.3349 ms | 75.2699 ms | — |
+| Cell triangle-configuration differences | 0 | 0 | 0 |
-| Full D3D wall time including upload, dispatch, and readback | 9.5323 ms summed | 136.3023 ms summed | — |
+| Protected/thin exception-cell configuration differences | 0 | 0 | 0 |
-| `World::generateTile` wall time, sequential sum | — | — | 15.8055 ms |
+| Weighted mean symmetric nearest-vertex distance | 0.000082887 cm | 0.000075131 cm | 0.000058929 cm |
-| CPU shape/sign export and upload preparation, sequential sum | — | — | 5.8304 ms |
+| Maximum symmetric nearest-vertex distance | 0.125 cm | 0.125 cm | 0.125 cm |
 
-The measured 0.3349 ms is compute-only field plus meshing across these seven probes. The seven-tile D3D wall total includes per-tile setup and readback and is not a viewer frame cost. The CPU API time is for the same seven keys, generated sequentially by the probe, not the full-view eight-worker total. On a 90 Hz display the full frame budget is 11.11 ms; these seven kernel dispatches account for 0.335 ms (3.0% of that budget), or fit within a 0.5 ms compute slice by themselves. This says nothing about scheduling all 1,266 LOD2+ tiles, viewer draw cost, or a steady-state frame because neither is integrated.
+There were 1,803,067 directed nearest-vertex comparisons. The metric is the symmetric mean/max of CPU-to-other and other-to-CPU nearest-vertex distances; it is not a continuous-surface Hausdorff bound. The largest per-tile coordinate-derived tolerance was 0.433 cm, and no tile exceeded its tolerance.
 
-The current CPU-only W8 clipmap benchmark took 0.691 s cold for 2,226 tiles in the baseline-source build. There is **no cold whole-view measurement with GPU far plus CPU near** yet, and no GPU chunk-crossing measurement in this step.
+The exception audit covered 127,401 cells and 353 route-adjacent thin-feature cells, with zero topology differences on either adapter. The input contained 29 route-protected-air samples, 89 thin-band samples, 89,528 seal-plane sign overrides, and no solid-overlay samples. Seal-plane overrides are exact CPU decisions; ordinary lattice signs are GPU decisions. Repeat RX runs produced identical output bits, cases, and indirect draw arguments on every tile.
 
-## Agreement and repeatability
+## GPU far-survival audit
 
-Across all 7 tiles (48,013 lattice samples), GPU-vs-CPU and WARP-vs-CPU had zero sign differences. The largest absolute field-value difference was `8.3923e-5` voxels on both adapters; per-tile mean absolute error was `5.98e-8` to `1.384e-6` voxels. Each GPU and WARP cell case was derived from the CPU’s exact sign mask, so there were zero cell case differences by construction. RX repeat runs had zero differing vertex-component bits, cases, or indirect-draw arguments on every tested tile.
+Run: `RunGpuFarViewer.ps1 -Mode audit -LogStem step3_gpu_far_audit_clip3_r2 -ClipRadius 3`. Strata boundaries passed: 7,200 boundary surfaces, with zero missing, extra, or hole counts. The full audit **failed on connectors**: across 10,976 far connectors, it sampled 178,351 points (157,407 visible and 20,944 with no source surface). It reported 156,345 missing visible connector samples and, over 710,047 connector rays, 10,168 missing, 11,724 extra, and 156,345 hole counts.
 
-The vertex metric is the symmetric mean/max of the two directed nearest-vertex distances (CPU-to-GPU and GPU-to-CPU); it is not a continuous-surface Hausdorff distance. Per tile:
+The GPU field mesh omits the CPU's `AppendFarConnectorShells` repair panels. The analytic connector field alone does not preserve those audited connector surfaces. The hybrid viewer runs, but this is a visible correctness gap in the GPU far mesh and the audit is not a pass. Use **H** to inspect the CPU far version where those repaired surfaces are required.
 
-| LOD | RX field max error / WARP max error (vox) | RX/CPU vertex mean / max (cm) | WARP/CPU vertex mean / max (cm) | RX/WARP vertex mean / max (cm) | Case diffs RX / WARP |
+## Viewer and cold timings
-| ---: | ---: | ---: | ---: | ---: | ---: |
-| 2 | 8.3923e-5 / 8.3923e-5 | 7.8943e-5 / 0.0029297 | 7.7662e-5 / 0.0029297 | 1.5237e-5 / 0.00097656 | 0 / 0 |
-| 3 | 9.5367e-6 / 9.5367e-6 | 7.0617e-5 / 0.00097656 | 7.5598e-5 / 0.00097656 | 6.8048e-6 / 0.00024414 | 0 / 0 |
-| 4 | 1.1444e-5 / 9.5367e-6 | 1.3605e-4 / 0.0097656 | 1.3279e-4 / 0.0097656 | 1.8673e-5 / 0.0019531 | 0 / 0 |
-| 5 | 1.5736e-5 / 1.5259e-5 | 1.8690e-4 / 0.0078125 | 2.4954e-4 / 0.0078125 | 1.4741e-4 / 0.0078125 | 0 / 0 |
-| 6 | 3.0518e-5 / 3.0518e-5 | 8.7221e-5 / 0.0078125 | 5.6992e-5 / 0.0078125 | 5.0486e-5 / 0.0039063 | 0 / 0 |
-| 7 | 3.0518e-5 / 3.0518e-5 | 1.0709e-4 / 0.015625 | 7.8328e-5 / 0.015625 | 7.7681e-5 / 0.015625 | 0 / 0 |
-| 8 | 2.2888e-5 / 2.2888e-5 | 1.5299e-4 / 0.03125 | 1.3583e-4 / 0.03125 | 5.6220e-5 / 0.015625 | 0 / 0 |
 
-Across tiles, the mean of the per-tile means was `0.0001171 cm` (about `1.17 µm`) for RX-vs-CPU and `0.0001152 cm` for WARP-vs-CPU. The largest per-tile maximum was `0.03125 cm` (0.3125 mm) for either CPU comparison. RX-vs-WARP mean-of-tile-means was `0.0000532 cm`, with a maximum of `0.015625 cm`. Their float outputs are not bit-identical: the field math differs slightly on the two implementations, while their sign masks and classified topology agree.
+The viewer draws GPU far vertex/index buffers directly, with CPU near and 152 CPU transition tiles. Cold runs were interleaved CPU A, GPU A, GPU B, CPU B using `RunGpuFarViewer.ps1` with `--quit-after-clipmap`, eight CPU workers, BelowNormal priority, and affinity `0xFF00`.
 
-No selected tile contained a protected-air sample or critical cell. Therefore these probes do not establish route-critical or thin-wall behavior; the route-critical sample count was zero, thin features were not audited, and the required “zero differences on protected/thin features” result remains unverified. The GPU far-survival audit against shell-repaired meshes was not run.
+| Cold clipmap | Run A | Run B | Mean |
+| --- | ---: | ---: | ---: |
+| All CPU far, 2,226 tiles | 2.206 s | 2.255 s | 2.231 s |
+| CPU near + GPU far, 1,114 GPU / 152 CPU transition tiles | 5.706 s | 5.770 s | 5.738 s |
 
-The separate 4,096-point shared relief-helper probe was bit-identical between C++ and RX/WARP in this run. That verifies the helper on those points; it is not an end-to-end relief-field comparison because the far shader does not apply relief.
+The hybrid took about **3.51 s longer** cold (2.57x the CPU run). GPU compute is fast, but CPU-side export/upload/tile preparation and one-tile-per-update scheduling dominate startup. The viewer's accumulated CPU input/upload/build work was 3.334 s and 3.402 s in the two GPU runs. CPU tile-worker time fell from 4.869 s to 1.671 s summed across workers, a 3.198 worker-second reduction that does not translate into a cold wall-time win because of that preparation cost.
+
+The viewer GPU timestamps summed to 32.620 ms and 32.753 ms across the 1,114 tiles. Since it schedules one tile per update, average field-plus-mesh compute was 0.0293–0.0294 ms per tile/frame and the slowest tile was 0.1216 ms. Against an 11.11 ms 90 Hz frame budget these are 0.26% average and 1.1% maximum. These are compute timestamps only, not full-frame GPU time. The separate agreement harness measured 10.885 ms of RX field dispatch and 46.456 ms of RX field-plus-meshing dispatch summed across its tiles; its setup and comparison workload differs from viewer streaming.
+
+The audit viewer peaked at 3.33 GB private memory; timed hybrid runs peaked at 3.28 GB private memory. The run scripts cap and record host private memory, but do not track a GPU-memory high-water mark. No `0xC0000142` occurred.
+
+The interactive viewer launched for the owner reached `[InitialClipmapReady]` in 5.579 s, submitted 1,114 GPU far tiles, and recorded 32.883 ms total GPU field-plus-mesh time (0.0295 ms mean, 0.1202 ms maximum per tile). It is left running for fly-through. Its latest process snapshot reports 3.37 GB private memory plus 2.52 GB dedicated/shared GPU process memory, about 5.88 GB combined. Priority is BelowNormal with affinity `0xFF00`.
 
 ## GTX 1060-class estimate
 
-Using the earlier first-port report’s nominal 61.4 TFLOPS RX 7900 XTX figure and the owner-provided 4.4 TFLOPS for the GTX 1060 class gives a peak-FP32 ratio of about 14:1. A direct shader-only scaling of the current seven-tile dispatch sum gives roughly **4.7 ms for field plus meshing across these seven probes** (about 0.67 ms per probe), of which the field-stage share is about 1.5 ms total. This is a rough extrapolation, not a measurement or a whole-view estimate. Pascal’s integer throughput, the field’s divergent candidate loops, clocks, and driver overhead make simple FP32 scaling especially uncertain.
+The RX 7900 XTX viewer measurement is about 0.0294 ms average and 0.122 ms maximum per GPU tile. Scaling only by the owner's nominal 4.4 TFLOPS GTX 1060-class fp32 figure against roughly 61.4 TFLOPS for the RX gives a **compute-only estimate** near 0.41 ms average and 1.7 ms maximum per tile, or about 0.46 s for the summed field-plus-mesh work across 1,114 tiles. Pascal's slower 32-bit integer multiply, divergent candidate loops, and driver overhead could push it higher; a practical estimate is several tenths of a millisecond to around 1 ms per average tile. This is not a GTX measurement. CPU input preparation remains the larger cold-start cost, so this estimate does not imply a whole-view speedup.
 
-## Missing deliverables
+## Captures and reproduction
 
-- The viewer does not yet draw the new tile buffers directly. The harness reads back GPU data for comparison.
+The same startup camera was captured in CPU and GPU modes: [CPU far](step3_cpu_far.png) and [GPU far](step3_gpu_far.png). Both are 1280x720 and appear visually indistinguishable from that camera; they do not validate remote connector survival. The GPU-mesh survival audit above is the relevant connector check.
-- There is no whole-view GPU-far/CPU-near cold timing or scheduled per-frame whole-far slice measurement.
-- Six fixed-view and four remote-view CPU-vs-GPU captures are not available.
-- The route-critical/thin-feature GPU test and far-survival audit against GPU meshes are not available.
-- GPU far shell repairs/connectors, the strata boundary list, and the relief term are not in the shader path; LOD0/1 detail relief has not been ported.
 
-## Reproduction
+Main logs: [whole-far agreement](step3_far_clip3_vertexcounts.stdout.log), [GPU survival audit](step3_gpu_far_audit_clip3_r2.stdout.log), and the four interleaved cold runs (`step3_cpu_timed_A/B.stdout.log`, `step3_gpu_timed_A/B.stdout.log`). Builds use the new `VoxelForgeGpuFarAgreement.exe` and `VoxelForgeGpuFarHybridViewer.exe` output names, `/MP8`, BelowNormal priority, and affinity `0xFF00`. All generated sources, binaries, logs, and captures remain under `Standalone/`. Nothing was committed, pushed, or stashed.
-
-The isolated build scripts are [BuildGpuFieldPort.ps1](BuildGpuFieldPort.ps1) and [BuildGpuFieldViewer.ps1](BuildGpuFieldViewer.ps1); the probe runner and final log are [RunGpuFieldPort.ps1](RunGpuFieldPort.ps1) and [far_field_lod2_8_fixed.stdout.log](far_field_lod2_8_fixed.stdout.log). Builds use `/MP8` and the host scripts and runner set BelowNormal priority with affinity `0xFF00`. The run reported peak process working set of 440.49 MiB and private memory of 611.45 MiB. The largest per-tile D3D buffer allocation was 3.135 MiB. No `0xC0000142` occurred. All build and result artifacts are under `Standalone\` with new output names; nothing was committed, pushed, or stashed.

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarAgreement.ps1

@@ -0,0 +1,42 @@
+[CmdletBinding()]
+param([string]$LogStem='whole_far_rx7900_warp',[switch]$WholeFar,[int]$ClipRadius=3)
+
+$ErrorActionPreference='Stop'
+if([IO.Path]::GetFileName($LogStem)-ne$LogStem){throw 'LogStem must be a filename stem.'}
+$HostProcess=Get-Process -Id $PID
+$HostProcess.PriorityClass='BelowNormal'
+$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$ExecutablePath=Join-Path $StandaloneRoot 'bin\VoxelForgeGpuFarAgreement.exe'
+$StdoutPath=Join-Path $PSScriptRoot ($LogStem+'.stdout.log')
+$StderrPath=Join-Path $PSScriptRoot ($LogStem+'.stderr.log')
+if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "GPU field runner not found: $ExecutablePath"}
+$Arguments=@("--clip-radius=$ClipRadius");if($WholeFar){$Arguments+='--whole-far'}
+
+for($Attempt=0;$Attempt -lt 2;$Attempt++)
+{
+    if($Attempt -gt 0)
+    {
+        Write-Host '0xC0000142 encountered; waiting 60 seconds before the single permitted retry.'
+        Start-Sleep -Seconds 60
+    }
+    Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
+    $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
+    $PeakWorkingSet=0L;$PeakPrivateMemory=0L
+    while(-not$Child.HasExited)
+    {
+        $Child.Refresh()
+        if($Child.HasExited){break}
+        $Child.PriorityClass='BelowNormal'
+        $Child.ProcessorAffinity=[IntPtr]0xFF00
+        $PeakWorkingSet=[Math]::Max($PeakWorkingSet,$Child.WorkingSet64)
+        $PeakPrivateMemory=[Math]::Max($PeakPrivateMemory,$Child.PrivateMemorySize64)
+        Start-Sleep -Milliseconds 100
+    }
+    $Child.WaitForExit();$Child.Refresh()
+    if($Child.ExitCode -eq -1073741502 -and $Attempt -eq 0){continue}
+    if($Child.ExitCode -eq -1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
+    Write-Output "RUN_EXIT pid=$($Child.Id) priority=BelowNormal affinity=0xFF00 peak_working_set_gb=$([math]::Round($PeakWorkingSet/1GB,2)) peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) exit_code=$($Child.ExitCode) stdout=$StdoutPath stderr=$StderrPath"
+    exit $Child.ExitCode
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarViewer.ps1

@@ -0,0 +1,71 @@
+[CmdletBinding()]
+param(
+    [ValidateSet('gpu','cpu','audit')][string]$Mode='gpu',
+    [string]$LogStem='gpu_far_viewer',
+    [int]$ClipRadius=3,
+    [string]$CaptureName=''
+)
+
+$ErrorActionPreference='Stop'
+if([IO.Path]::GetFileName($LogStem)-ne$LogStem){throw 'LogStem must be a filename stem.'}
+$HostProcess=Get-Process -Id $PID
+$HostProcess.PriorityClass='BelowNormal'
+$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$ExecutablePath=Join-Path $StandaloneRoot 'bin\VoxelForgeGpuFarHybridViewer.exe'
+$StdoutPath=Join-Path $PSScriptRoot ($LogStem+'.stdout.log')
+$StderrPath=Join-Path $PSScriptRoot ($LogStem+'.stderr.log')
+if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "GPU far viewer not found: $ExecutablePath"}
+$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($CaptureName)
+{
+    if([IO.Path]::GetFileName($CaptureName)-ne$CaptureName-or[IO.Path]::GetExtension($CaptureName)-ne'.png')
+    {throw 'CaptureName must be a PNG filename.'}
+    $Arguments+='--capture-after-clipmap'
+    $Arguments+=("--capture=Research/OwnerShapeWorld/gpu_field_port/$CaptureName")
+}
+$Command="`"$ExecutablePath`" $($Arguments -join ' ')"
+Write-Output "VIEWER_COMMAND mode=$Mode $Command"
+
+for($Attempt=0;$Attempt -lt 2;$Attempt++)
+{
+    if($Attempt-gt 0)
+    {
+        Write-Host '0xC0000142 encountered; waiting 60 seconds before the single permitted retry.'
+        Start-Sleep -Seconds 60
+    }
+    Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
+    $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
+    $PeakWorkingSet=0L;$PeakPrivateMemory=0L;$OverMemoryLimit=$false
+    while(-not$Child.HasExited)
+    {
+        $Child.Refresh()
+        if($Child.HasExited){break}
+        $Child.PriorityClass='BelowNormal'
+        $Child.ProcessorAffinity=[IntPtr]0xFF00
+        $PeakWorkingSet=[Math]::Max($PeakWorkingSet,$Child.WorkingSet64)
+        $PeakPrivateMemory=[Math]::Max($PeakPrivateMemory,$Child.PrivateMemorySize64)
+        if($Child.PrivateMemorySize64-ge8GB)
+        {
+            $OverMemoryLimit=$true
+            $Child.Kill();$Child.WaitForExit();break
+        }
+        Start-Sleep -Milliseconds 100
+    }
+    $Child.WaitForExit();$Child.Refresh()
+    if($Child.ExitCode-eq-1073741502-and$Attempt-eq0){continue}
+    if($Child.ExitCode-eq-1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
+    if($OverMemoryLimit){throw "Viewer stopped at the 8 GB per-process memory limit; stdout=$StdoutPath stderr=$StderrPath"}
+    $LogText=Get-Content -LiteralPath $StdoutPath -Raw
+    $Ready=$LogText.Contains('[InitialClipmapReady]')
+    $AuditComplete=($Mode-ne'audit')-or$LogText.Contains('[GpuFarSurvivalAudit]')
+    $AuditState=if($Mode-eq'audit'){if($AuditComplete){'reported'}else{'missing'}}else{'not-requested'}
+    Write-Output "RUN_EXIT pid=$($Child.Id) priority=BelowNormal affinity=0xFF00 peak_working_set_gb=$([math]::Round($PeakWorkingSet/1GB,2)) peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) exit_code=$($Child.ExitCode) clipmap_ready=$Ready audit=$AuditState stdout=$StdoutPath stderr=$StderrPath"
+    if(-not$Ready){exit 4}
+    if(-not$AuditComplete){exit 5}
+    exit $Child.ExitCode
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/StartGpuFarViewer.ps1

@@ -0,0 +1,16 @@
+[CmdletBinding()]
+param([int]$ClipRadius=3)
+
+$ErrorActionPreference='Stop'
+$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$ExecutablePath=Join-Path $StandaloneRoot 'bin\VoxelForgeGpuFarHybridViewer.exe'
+$StdoutPath=Join-Path $PSScriptRoot 'owner_hybrid_viewer.stdout.log'
+$StderrPath=Join-Path $PSScriptRoot 'owner_hybrid_viewer.stderr.log'
+if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "GPU far viewer not found: $ExecutablePath"}
+$Arguments=@('--gpu-far',"--clip-radius=$ClipRadius",'--workers=8')
+Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
+$Viewer=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
+    -WindowStyle Normal -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
+$Viewer.PriorityClass='BelowNormal'
+$Viewer.ProcessorAffinity=[IntPtr]0xFF00
+Write-Output "STARTED pid=$($Viewer.Id) priority=BelowNormal affinity=0xFF00 command=`"$ExecutablePath`" $($Arguments -join ' ') stdout=$StdoutPath stderr=$StderrPath"

Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl

@@ -130,9 +130,10 @@ struct FarFieldSampleRange { uint offset; uint count; };
 StructuredBuffer<VF_FieldPrimitive> FarShapes : register(t10);
 StructuredBuffer<FarFieldSampleRange> FarSampleShapeRanges : register(t11);
 StructuredBuffer<uint> FarSampleShapeIds : register(t12);
-StructuredBuffer<uint> FarRockAirSigns : register(t13);
+StructuredBuffer<uint> FarProtectedAir : register(t13);
-StructuredBuffer<uint> FarProtectedAir : register(t14);
+StructuredBuffer<uint> FarProtectedSolid : register(t14);
 RWStructuredBuffer<float> FarFieldDensity : register(u7);
+StructuredBuffer<float> GpuSignDensity : register(t15);
 
 cbuffer FarFieldConstants : register(b1)
 {
@@ -186,31 +187,22 @@ void EvaluateFarFieldCS(uint3 dispatchId : SV_DispatchThreadID)
     const float base = min(position.z - FarStrataLo, FarStrataHi - position.z);
     float rockDistance = VF_FieldSmoothMin(base, max(cave, overlay), FarStrataBlend);
     const uint protectedAir = (FarProtectedAir[index >> 5] >> (index & 31)) & 1;
-    if (protectedAir != 0) rockDistance = min(rockDistance, -0.0625f);
+    const uint protectedSolid = (FarProtectedSolid[index >> 5] >> (index & 31)) & 1;
     const float rawDensity = -rockDistance;
-    const uint air = (FarRockAirSigns[index >> 5] >> (index & 31)) & 1;
     const float magnitude = max(abs(rawDensity), 0.05f);
-    FarFieldDensity[index] = air != 0 ? magnitude : -magnitude;
+    float density = rockDistance >= 0.0f ? -magnitude : magnitude;
+    if (protectedAir != 0) density = magnitude;
+    else if (protectedSolid != 0) density = -magnitude;
+    FarFieldDensity[index] = density;
 }
 
 groupshared uint ScanValues[256];
 
-[numthreads(256,1,1)]
+void ClassifyAndScan(uint index, uint3 groupId, uint3 groupThreadId, uint caseIndex)
-void ClassifyCS(uint3 dispatchId : SV_DispatchThreadID,
-                uint3 groupId : SV_GroupID,
-                uint3 groupThreadId : SV_GroupThreadID)
 {
-    const uint index = dispatchId.x;
     uint vertexCount = 0;
-    uint caseIndex = 0;
     if (index < CellCount)
     {
-        const uint3 c = uint3(index % CellsPerAxis,
-                              (index / CellsPerAxis) % CellsPerAxis,
-                              index / (CellsPerAxis * CellsPerAxis));
-        [unroll]
-        for (uint corner = 0; corner < 8; ++corner)
-            caseIndex |= LatticeSign(Signs, int3(c) + CornerOffset[corner]) << corner;
         [unroll]
         for (uint slot = 0; slot < 15; ++slot)
         {
@@ -235,6 +227,49 @@ void ClassifyCS(uint3 dispatchId : SV_DispatchThreadID,
         GroupVertexCounts[groupId.x] = ScanValues[255];
 }
 
+[numthreads(256,1,1)]
+void ClassifyCS(uint3 dispatchId : SV_DispatchThreadID,
+                uint3 groupId : SV_GroupID,
+                uint3 groupThreadId : SV_GroupThreadID)
+{
+    const uint index = dispatchId.x;
+    uint caseIndex = 0;
+    if (index < CellCount)
+    {
+        const uint3 c = uint3(index % CellsPerAxis,
+                              (index / CellsPerAxis) % CellsPerAxis,
+                              index / (CellsPerAxis * CellsPerAxis));
+        [unroll]
+        for (uint corner = 0; corner < 8; ++corner)
+            caseIndex |= LatticeSign(Signs, int3(c) + CornerOffset[corner]) << corner;
+    }
+    ClassifyAndScan(index, groupId, groupThreadId, caseIndex);
+}
+
+uint GpuLatticeSign(int3 p)
+{
+    return GpuSignDensity[SampleIndex(p)] >= 0.0f ? 1 : 0;
+}
+
+[numthreads(256,1,1)]
+void ClassifyFarCS(uint3 dispatchId : SV_DispatchThreadID,
+                   uint3 groupId : SV_GroupID,
+                   uint3 groupThreadId : SV_GroupThreadID)
+{
+    const uint index = dispatchId.x;
+    uint caseIndex = 0;
+    if (index < CellCount)
+    {
+        const uint3 c = uint3(index % CellsPerAxis,
+                              (index / CellsPerAxis) % CellsPerAxis,
+                              index / (CellsPerAxis * CellsPerAxis));
+        [unroll]
+        for (uint corner = 0; corner < 8; ++corner)
+            caseIndex |= GpuLatticeSign(int3(c) + CornerOffset[corner]) << corner;
+    }
+    ClassifyAndScan(index, groupId, groupThreadId, caseIndex);
+}
+
 StructuredBuffer<uint> InputGroupVertexCounts : register(t8);
 RWStructuredBuffer<uint> GroupOffsets : register(u3);
 RWByteAddressBuffer DrawArgs : register(u4);

Standalone/Source/StandaloneViewer.cpp

@@ -2,6 +2,7 @@
 #include "StreamingWorld.h"
 #include "StandaloneCoreScheduler.h"
 #include "../Core/src/StandaloneSnapshotAdapter.h"
+#include "../Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h"
 #include "VoxelDensityProfile.h"
 
 #ifdef TEXT
@@ -23,6 +24,7 @@
 #include <complex>
 #include <cstdio>
 #include <cstdlib>
+#include <deque>
 #include <filesystem>
 #include <fstream>
 #include <iomanip>
@@ -1345,6 +1347,16 @@ namespace
         std::chrono::steady_clock::time_point StartTime{};
         bool bReported=false;
     };
+    struct FGpuFarTile
+    {
+        ID3D11Buffer* VertexBuffer=nullptr;
+        ID3D11Buffer* IndexBuffer=nullptr;
+        ID3D11Buffer* IndirectArgs=nullptr;
+        ID3D11Query* CompletionEvent=nullptr;
+        uint64 VertexCapacityBytes=0;
+        double GpuDispatchMs=0.0;
+        bool bGpuComplete=false;
+    };
 
     uint64 HashStandaloneFile(const std::string& Path)
     {
@@ -1426,6 +1438,10 @@ namespace
         std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> DesiredTileSet;
         std::unordered_map<FStandaloneBatchKey,std::vector<FVoxelTileKey>,FStandaloneBatchKeyHash> BatchMembers;
         std::unordered_map<FStandaloneBatchKey,FStreamingGpuBatch,FStandaloneBatchKeyHash> StreamingBatches;
+        std::unordered_map<FVoxelTileKey,FGpuFarTile,FStandaloneTileKeyHash> GpuFarTiles;
+        std::deque<FVoxelTileKey> GpuFarQueue;
+        std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> GpuFarQueued;
+        std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> GpuFarEligibleTiles;
         std::unordered_set<FStandaloneBatchKey,FStandaloneBatchKeyHash> DirtyBatchSet;
         std::unordered_map<uint64,FCrossingWork> CrossingWork;
         std::vector<FVoxelTileKey> DesiredTiles;
@@ -1448,6 +1464,16 @@ namespace
         bool bShapeAuditOnly=false;
         bool bShapeCompare=false;
         bool bFarShapeStagePending=false;
+        bool bGpuFarEnabled=false;
+        bool bGpuFarAuditEnabled=false;
+        bool bGpuFarAuditReported=false;
+        std::unique_ptr<FFarSurvivalAudit> GpuFarAudit;
+        bool bGpuFarFirstTileReported=false;
+        std::filesystem::path GpuFarShaderPath;
+        double GpuFarEnqueueWallSeconds=0.0;
+        uint64 GpuFarTilesSubmitted=0;
+        double GpuFarGpuDispatchMs=0.0;
+        double GpuFarGpuDispatchMaxMs=0.0;
         VoxelForgeCore::DetailVariant ShapeDetailVariant=VoxelForgeCore::DetailVariant::coherent;
         bool bWireframe=false;
         bool bCaptured=false;
@@ -1513,6 +1539,17 @@ namespace
                 if(Pair.second.IndexBuffer)Pair.second.IndexBuffer->Release();
             }
             StreamingBatches.clear();BatchMembers.clear();ResidentTiles.clear();
+            for(auto& Pair:GpuFarTiles)
+            {
+                if(Pair.second.VertexBuffer)Pair.second.VertexBuffer->Release();
+                if(Pair.second.IndexBuffer)Pair.second.IndexBuffer->Release();
+                if(Pair.second.IndirectArgs)Pair.second.IndirectArgs->Release();
+                if(Pair.second.CompletionEvent)Pair.second.CompletionEvent->Release();
+            }
+            GpuFarTiles.clear();GpuFarQueue.clear();GpuFarQueued.clear();GpuFarEligibleTiles.clear();
+            GpuFarAudit.reset();bGpuFarAuditReported=false;
+            GpuFarTilesSubmitted=0;GpuFarEnqueueWallSeconds=0.0;GpuFarGpuDispatchMs=0.0;
+            GpuFarGpuDispatchMaxMs=0.0;bGpuFarFirstTileReported=false;
             DirtyBatchSet.clear();DesiredTiles.clear();DesiredTileSet.clear();CrossingWork.clear();
             DeferredLeaverTiles.clear();
             for(int32 I=0;I<=MaxLevels;++I)
@@ -1668,7 +1705,7 @@ namespace
 cbuffer C:register(b0){float4x4 VP;float3 Camera;float FogStart;float3 Forward;float FogEnd;float3 LightDir;float Ambient;float4 Pad;}
 struct I{float3 p:POSITION;float3 n:NORMAL;};
 struct O{float4 p:SV_POSITION;float3 world:TEXCOORD0;float3 n:NORMAL;};
-O VS(I i){O o;o.p=mul(float4(i.p,1),VP);o.world=i.p;o.n=normalize(i.n);return o;}
+ O VS(I i){O o;float3 p=i.p*Pad.x;o.p=mul(float4(p,1),VP);o.world=p;o.n=normalize(i.n);return o;}
 float4 PS(O i):SV_TARGET{
  float3 N=normalize(i.n);float3 V=normalize(Camera-i.world);float3 L=normalize(-LightDir);
  float Diffuse=saturate(dot(N,L));float Head=saturate(dot(N,V));
@@ -1974,6 +2011,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 (!bShapeWorld&&MeshCache.IsLoaded())?&MeshCache:nullptr,
                 bFarShapeStagePending?World->FullResClipLevels-1:-1,
                 false,false,false,false,ShapeWorld.get());
+            if(bGpuFarEnabled&&ShapeWorld&&ShapeWorld->hasFarWorld())
+            {
+                ScheduleGpuFarTiles(DesiredTiles);
+                if(bGpuFarAuditEnabled)
+                {
+                    std::vector<FVoxelTileKey> AuditKeys;AuditKeys.reserve(GpuFarEligibleTiles.size());
+                    for(const FVoxelTileKey& Key:DesiredTiles)
+                        if(GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key);
+                    GpuFarAudit=std::make_unique<FFarSurvivalAudit>(*World,*ShapeWorld,AuditKeys,false);
+                }
+            }
             if(bShapePipelineActive)Streamer->SetDesiredTileFilter(InitialNeighborhoodTiles);
             else if(CaptureTileLimit>0&&CapturePath.size())
                 Streamer->SetDesiredTileFilter(DesiredTiles);
@@ -2025,6 +2073,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             Streamer.reset();
             Streamer=std::make_unique<FStandaloneStreamingGenerator>(*World,WorkerCount,false,nullptr,-1,
                 false,false,false,false,ShapeWorld.get());
+            if(bGpuFarEnabled)
+            {
+                ScheduleGpuFarTiles(DesiredTiles);
+                if(bGpuFarAuditEnabled&&World)
+                {
+                    std::vector<FVoxelTileKey> AuditKeys;AuditKeys.reserve(GpuFarEligibleTiles.size());
+                    for(const FVoxelTileKey& Key:DesiredTiles)
+                        if(GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key);
+                    GpuFarAudit=std::make_unique<FFarSurvivalAudit>(*World,*ShapeWorld,AuditKeys,false);
+                }
+            }
             const FVector CameraVoxels(Position.X*4.0,Position.Y*4.0,Position.Z*4.0);
             ++CrossingIndex;
             const FStandaloneStreamingUpdate Update=Streamer->Start(CameraVoxels,ResidentKeys,CrossingIndex);
@@ -2037,6 +2096,145 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             std::fflush(stdout);
         }
 
+        bool IsRenderedTileReady(const FVoxelTileKey& Key)
+        {
+            if(bGpuFarEnabled&&ShapeWorld&&ShapeWorld->hasFarWorld()
+                &&GpuFarEligibleTiles.contains(Key))
+            {
+                auto It=GpuFarTiles.find(Key);
+                if(It==GpuFarTiles.end())return false;
+                FGpuFarTile& Tile=It->second;
+                if(Tile.bGpuComplete)return true;
+                if(Tile.CompletionEvent&&Context
+                    &&Context->GetData(Tile.CompletionEvent,nullptr,0,D3D11_ASYNC_GETDATA_DONOTFLUSH)==S_OK)
+                {
+                    Tile.bGpuComplete=true;
+                    Tile.CompletionEvent->Release();Tile.CompletionEvent=nullptr;
+                }
+                return Tile.bGpuComplete;
+            }
+            return ResidentTiles.contains(Key);
+        }
+
+        void ScheduleGpuFarTiles(const std::vector<FVoxelTileKey>& Keys)
+        {
+            GpuFarQueue.clear();GpuFarQueued.clear();GpuFarEligibleTiles.clear();
+            if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld())return;
+            std::vector<FVoxelTileKey> CpuTiles;CpuTiles.reserve(Keys.size());
+            for(const FVoxelTileKey& Key:Keys)
+            {
+                if(Key.Level<World->FullResClipLevels){CpuTiles.push_back(Key);continue;}
+                const VoxelForgeCore::TileKey CoreKey{{Key.Coord.X,Key.Coord.Y,Key.Coord.Z},Key.Level};
+                if(ShapeWorld->isFarFieldTileEligible(CoreKey))
+                {
+                    GpuFarEligibleTiles.insert(Key);
+                    if(!GpuFarTiles.contains(Key))
+                    {GpuFarQueue.push_back(Key);GpuFarQueued.insert(Key);}
+                }
+                else CpuTiles.push_back(Key);
+            }
+            if(Streamer)Streamer->SetDesiredTileFilter(CpuTiles);
+            const std::size_t CpuTransitionFallbacks=static_cast<std::size_t>(std::count_if(
+                Keys.begin(),Keys.end(),[&](const FVoxelTileKey& Key)
+                    {return Key.Level>=World->FullResClipLevels&&!GpuFarEligibleTiles.contains(Key);}));
+            std::printf("[GpuFarEligibility] eligible_gpu_tiles=%zu cpu_transition_fallback_tiles=%zu\n",
+                GpuFarEligibleTiles.size(),CpuTransitionFallbacks);
+        }
+
+        void PumpGpuFarGeneration()
+        {
+            if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld()||!Device||!Context
+                ||GpuFarQueue.empty())return;
+            const auto Start=std::chrono::steady_clock::now();
+            int32 Built=0;
+            do
+            {
+                const FVoxelTileKey Key=GpuFarQueue.front();GpuFarQueue.pop_front();GpuFarQueued.erase(Key);
+                if(!DesiredTileSet.contains(Key)||GpuFarTiles.contains(Key))continue;
+                FGpuFarTile Tile;
+                VoxelForgeGpuFarViewer::CpuMesh AuditMesh;
+                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};
+                if(!VoxelForgeGpuFarViewer::BuildFarTile(Device,Context,*ShapeWorld,CoreKey,
+                    GpuFarShaderPath,&Tile.VertexBuffer,&Tile.IndexBuffer,&Tile.IndirectArgs,
+                    &Tile.CompletionEvent,
+                    (bGpuFarAuditEnabled&&Key.Level>=6)?&AuditMesh:nullptr,
+                    &EnqueueMs,&GpuDispatchMs,Error))
+                {
+                    std::fprintf(stderr,"[GpuFarTileError] key=(%d,%d,%d,lod%d): %s\n",
+                        Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,Error.c_str());
+                    bGpuFarEnabled=false;
+                    break;
+                }
+                if(bGpuFarAuditEnabled&&GpuFarAudit&&Key.Level>=6)
+                {
+                    FStandaloneCpuTile AuditTile;AuditTile.Key=Key;
+                    AuditTile.SurfaceVertexCount=static_cast<uint32>(AuditMesh.vertices.size());
+                    AuditTile.Vertices.reserve(AuditMesh.vertices.size());
+                    for(const auto& Vertex:AuditMesh.vertices)
+                        AuditTile.Vertices.push_back({
+                            {Vertex.position[0]*0.01f,Vertex.position[1]*0.01f,Vertex.position[2]*0.01f},
+                            {Vertex.normal[0],Vertex.normal[1],Vertex.normal[2]}});
+                    AuditTile.Indices=std::move(AuditMesh.indices);
+                    GpuFarAudit->AuditTile(AuditTile);
+                }
+                const uint64 Cells=static_cast<uint64>(World->CoarseTileCells);
+                Tile.VertexCapacityBytes=Cells*Cells*Cells*15u*32u;
+                Tile.GpuDispatchMs=GpuDispatchMs;
+                GpuFarEnqueueWallSeconds+=EnqueueMs/1000.0;
+                GpuFarGpuDispatchMs+=GpuDispatchMs;
+                GpuFarGpuDispatchMaxMs=std::max(GpuFarGpuDispatchMaxMs,GpuDispatchMs);
+                ++GpuFarTilesSubmitted;++Built;
+                GpuFarTiles.emplace(Key,Tile);
+                if(!bGpuFarFirstTileReported)
+                {
+                    bGpuFarFirstTileReported=true;
+                    std::printf("[GpuFarFirstTile] key=(%d,%d,%d,lod%d) cpu_export_upload_enqueue_ms=%.3f vertex_capacity_mb=%.3f mode=direct-indirect-draw\n",
+                        Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,EnqueueMs,
+                        Tile.VertexCapacityBytes/(1024.0*1024.0));
+                    std::fflush(stdout);
+                }
+            }
+            while(Built<1&& !GpuFarQueue.empty()
+                &&std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-Start).count()<1.0);
+        }
+
+        void ToggleFarRenderer()
+        {
+            if(!bShapeWorld||!ShapeWorld||!ShapeWorld->hasFarWorld())
+            {
+                bGpuFarEnabled=!bGpuFarEnabled;
+                std::printf("[FarRenderer] GPU far %s; it will take effect when far shapes are ready.\n",
+                    bGpuFarEnabled?"enabled":"disabled");
+                return;
+            }
+            bGpuFarEnabled=!bGpuFarEnabled;
+            if(bGpuFarEnabled)
+            {
+                ScheduleGpuFarTiles(DesiredTiles);
+                std::vector<FVoxelTileKey> CpuFar;
+                for(const FVoxelTileKey& Key:DesiredTiles)
+                    if(GpuFarEligibleTiles.contains(Key)&&ResidentTiles.contains(Key))CpuFar.push_back(Key);
+                for(const FVoxelTileKey& Key:CpuFar)RemoveResidentTile(Key);
+            }
+            else
+            {
+                GpuFarQueue.clear();GpuFarQueued.clear();GpuFarEligibleTiles.clear();
+                Streamer->SetDesiredTileFilter(DesiredTiles);
+            }
+            std::vector<FVoxelTileKey> ResidentKeys;ResidentKeys.reserve(ResidentTiles.size());
+            for(const auto& Pair:ResidentTiles)ResidentKeys.push_back(Pair.first);
+            const FVector CameraVoxels(Position.X*4.0,Position.Y*4.0,Position.Z*4.0);
+            ++CrossingIndex;
+            const FStandaloneStreamingUpdate Update=Streamer->UpdateDesired(
+                CameraVoxels,ResidentKeys,CrossingIndex);
+            FCrossingWork& Work=CrossingWork[CrossingIndex];
+            Work.SubmittedJobs=Update.NewJobs;Work.PendingJobs=Update.NewJobs;
+            Work.StartTime=std::chrono::steady_clock::now();
+            std::printf("[FarRenderer] mode=%s cpu_far_jobs=%d gpu_far_ready=%zu gpu_far_queued=%zu (H toggles)\n",
+                bGpuFarEnabled?"GPU":"CPU",Update.NewJobs,GpuFarTiles.size(),GpuFarQueue.size());
+        }
+
         void TryCompleteShapePipeline()
         {
             if(!bShapePipelineActive||InitialNeighborhoodTiles.empty())return;
@@ -2222,6 +2420,16 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             }
             DesiredTiles=NewDesired.Tiles;
             DesiredTileSet=std::move(NewSet);
+            for(auto It=GpuFarTiles.begin();It!=GpuFarTiles.end();)
+            {
+                if(DesiredTileSet.contains(It->first)){++It;continue;}
+                if(It->second.VertexBuffer)It->second.VertexBuffer->Release();
+                if(It->second.IndexBuffer)It->second.IndexBuffer->Release();
+                if(It->second.IndirectArgs)It->second.IndirectArgs->Release();
+                if(It->second.CompletionEvent)It->second.CompletionEvent->Release();
+                It=GpuFarTiles.erase(It);
+            }
+            ScheduleGpuFarTiles(DesiredTiles);
             RefreshResidentSkirtMasks();
             for(int32 I=0;I<=MaxLevels;++I)DesiredByLevel[I]=0;
             for(const FVoxelTileKey& Tile:DesiredTiles)++DesiredByLevel[Tile.Level];
@@ -2489,7 +2697,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             {
                 bool bDesiredResident=true;
                 for(const FVoxelTileKey& Key:DesiredTiles)
-                    if(!ResidentTiles.contains(Key)){bDesiredResident=false;break;}
+                    if(!IsRenderedTileReady(Key)){bDesiredResident=false;break;}
                 if(bDesiredResident)
                 {
                     std::vector<FVoxelTileKey> ReadyToDrop;
@@ -2531,7 +2739,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     const uint64 CacheHits=Streamer?Streamer->GetCacheHits():0;
                     std::size_t DesiredResidentCount=0;
                     for(const FVoxelTileKey& Key:DesiredTiles)
-                        if(ResidentTiles.contains(Key))++DesiredResidentCount;
+                        if(IsRenderedTileReady(Key))++DesiredResidentCount;
                     std::sort(NearestReadyMilliseconds.begin(),NearestReadyMilliseconds.end());
                     const auto Percentile=[&](std::size_t Percent)
                     {
@@ -2568,7 +2776,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             }
             size_t DesiredResidentCount=0;
             for(const FVoxelTileKey& Key:DesiredTiles)
-                if(ResidentTiles.contains(Key))++DesiredResidentCount;
+                if(IsRenderedTileReady(Key))++DesiredResidentCount;
             if(!bInitialClipmapReported&&DesiredResidentCount==DesiredTiles.size()&&DirtyBatchSet.empty())
             {
                 bInitialClipmapReported=true;
@@ -2576,10 +2784,28 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 std::printf("[InitialClipmapReady] %.3f s after cold process start; %zu tiles, %llu vertices, %llu indices, %zu GPU batches.\n",
                     Elapsed,DesiredResidentCount,static_cast<unsigned long long>(TotalVertices),
                     static_cast<unsigned long long>(TotalIndices),StreamingBatches.size());
+                if(bGpuFarEnabled)
+                {
+                    const double AverageGpuMs=GpuFarTilesSubmitted
+                        ?GpuFarGpuDispatchMs/static_cast<double>(GpuFarTilesSubmitted):0.0;
+                    std::printf("[GpuFarTiming] submitted_tiles=%llu cpu_export_upload_enqueue_wall_ms=%.3f gpu_timestamp_total_ms=%.3f gpu_timestamp_mean_per_one_tile_frame_ms=%.6f gpu_timestamp_max_tile_ms=%.6f frame_budget_90hz_ms=11.111\n",
+                        static_cast<unsigned long long>(GpuFarTilesSubmitted),
+                        GpuFarEnqueueWallSeconds*1000.0,GpuFarGpuDispatchMs,AverageGpuMs,GpuFarGpuDispatchMaxMs);
+                }
                 ReportTileTimings();std::fflush(stdout);
                 if(bQuitAfterClipmap&&CapturePath.empty())PostMessageA(Window,WM_CLOSE,0,0);
             }
-            bGenerationRunning=DesiredResidentCount!=DesiredTiles.size()||!DirtyBatchSet.empty();
+            if(bInitialClipmapReported&&bGpuFarAuditEnabled&&GpuFarAudit&&!bGpuFarAuditReported)
+            {
+                GpuFarAudit->Finalize();GpuFarAudit->Print();
+                bGpuFarAuditReported=true;
+                std::printf("[GpuFarSurvivalAudit] mesh_source=GPU field_readback=required status=%s\n",
+                    GpuFarAudit->Passing()?"PASS":"FAIL");
+                std::fflush(stdout);
+                if(bQuitAfterClipmap&&CapturePath.empty())PostMessageA(Window,WM_CLOSE,0,0);
+            }
+            bGenerationRunning=DesiredResidentCount!=DesiredTiles.size()||!DirtyBatchSet.empty()
+                ||(bGpuFarEnabled&&!GpuFarQueue.empty());
         }
 
         void CaptureRenderedFrame()
@@ -2608,7 +2834,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             uint64 VisibleIndices=0;for(const auto& Pair:StreamingBatches)if(bShowLevel[Pair.first.Level])VisibleIndices+=Pair.second.IndexCount;
             std::ostringstream S;S<<"VoxelForge standalone | seed "<<Seed<<" | "<<(bShapeWorld?"shape-interval owner world":(bRoundTunnels?"round/winding research tunnels":"base world"))
                 <<(bRoundTunnels?(bRoundTunnelEnhanced?" (detailed walls)":" (legacy walls)"):"")
-                <<" | "<<VisibleIndices/3<<" tris | "<<FrameMs<<" ms | "<<(bGenerationRunning?"generating":"ready")
+                <<" | "<<VisibleIndices/3<<" tris | "<<FrameMs<<" ms | far "<<(bGpuFarEnabled?"GPU":"CPU")
+                <<" | "<<(bGenerationRunning?"generating":"ready")
                 <<" | pos "<<std::fixed<<std::setprecision(1)<<Position.X<<", "<<Position.Y<<", "<<Position.Z<<" m";
             SetWindowTextA(Window,S.str().c_str());
         }
@@ -2616,6 +2843,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         void Draw()
         {
             PumpGeneration();
+            PumpGpuFarGeneration();
             const auto Now=std::chrono::steady_clock::now();
             if(FlightTargetCrossings>0&&bFullDetailReported&&!bFlightStarted)
             {
@@ -2658,6 +2886,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             C.ViewProjection=Transpose(Combined);C.Camera[0]=Position.X;C.Camera[1]=Position.Y;C.Camera[2]=Position.Z;
             C.Forward[0]=F.X;C.Forward[1]=F.Y;C.Forward[2]=F.Z;C.LightDirection[0]=-0.5f;C.LightDirection[1]=-0.72f;C.LightDirection[2]=-0.48f;
             C.FogStart=260.0f;C.FogEnd=2600.0f;C.Ambient=0.36f;
+            C.Padding[0]=1.0f;
             Context->UpdateSubresource(Constants,0,nullptr,&C,0,0);Context->VSSetConstantBuffers(0,1,&Constants);Context->PSSetConstantBuffers(0,1,&Constants);
             Context->VSSetShader(VertexShader,nullptr,0);Context->PSSetShader(PixelShader,nullptr,0);
             for(const auto& Pair:StreamingBatches)
@@ -2667,6 +2896,21 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 const UINT Stride=sizeof(FGpuVertex),Offset=0;Context->IASetVertexBuffers(0,1,&L.VertexBuffer,&Stride,&Offset);
                 Context->IASetIndexBuffer(L.IndexBuffer,DXGI_FORMAT_R32_UINT,0);Context->DrawIndexed(L.IndexCount,0,0);
             }
+            if(bGpuFarEnabled&&ShapeWorld&&ShapeWorld->hasFarWorld())
+            {
+                C.Padding[0]=0.01f;
+                Context->UpdateSubresource(Constants,0,nullptr,&C,0,0);
+                Context->VSSetConstantBuffers(0,1,&Constants);Context->PSSetConstantBuffers(0,1,&Constants);
+                for(const auto& Pair:GpuFarTiles)
+                {
+                    const FGpuFarTile& Tile=Pair.second;
+                    if(!bShowLevel[Pair.first.Level]||!Tile.VertexBuffer||!Tile.IndexBuffer||!Tile.IndirectArgs)continue;
+                    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,0);
+                }
+            }
             CaptureRenderedFrame();
             SwapChain->Present(1,0);UpdateTitle();
         }
@@ -2693,6 +2937,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 std::printf("Round-tunnel comparison: %s wall field.\n",
                     GApp->bRoundTunnelEnhanced?"detailed/smoothed":"legacy smooth tube");
             }
+            if(!bWasDown&&W=='H')GApp->ToggleFarRenderer();
             if(!bWasDown&&(W=='N'||W=='P'))GApp->RequestSeed(W=='N'?1:-1);
             if(!bWasDown&&W>='1'&&W<='8'){const int32 I=static_cast<int32>(W-'1');GApp->bShowLevel[I]=!GApp->bShowLevel[I];}
             if(!bWasDown&&W=='9')GApp->bShowLevel[8]=!GApp->bShowLevel[8];
@@ -4425,8 +4670,15 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
 
 int main(int Argc,char** Argv)
 {
+    SetPriorityClass(GetCurrentProcess(),BELOW_NORMAL_PRIORITY_CLASS);
+    SetProcessAffinityMask(GetCurrentProcess(),static_cast<DWORD_PTR>(0xFF00));
     const auto ColdProcessStart=std::chrono::steady_clock::now();
     FApp App;App.DataPath=DefaultDataPath().string();App.ColdProcessStart=ColdProcessStart;
+    wchar_t ExecutablePath[MAX_PATH]{};
+    const DWORD ExecutablePathLength=GetModuleFileNameW(nullptr,ExecutablePath,MAX_PATH);
+    if(ExecutablePathLength>0&&ExecutablePathLength<MAX_PATH)
+        App.GpuFarShaderPath=std::filesystem::path(ExecutablePath).parent_path().parent_path()
+            /"Research"/"OwnerShapeWorld"/"gpu_first_port"/"VoxelForgeGpuMesher.hlsl";
     std::string JoinedCommandLine;
     for(int32 I=1;I<Argc;++I)
     {if(!JoinedCommandLine.empty())JoinedCommandLine.push_back(' ');JoinedCommandLine+=Argv[I];}
@@ -4464,6 +4716,10 @@ int main(int Argc,char** Argv)
         else if(Arg=="--parity-only")App.bParityOnly=true;
         else if(Arg=="--density-hash")App.bDensityHash=true;
         else if(Arg=="--shape-world")App.bShapeWorld=true;
+        else if(Arg=="--gpu-far")App.bGpuFarEnabled=true,App.bShapeWorld=true,App.bForceGenerate=true;
+        else if(Arg=="--cpu-far")App.bGpuFarEnabled=false,App.bShapeWorld=true,App.bForceGenerate=true;
+        else if(Arg=="--gpu-far-audit")
+        {App.bGpuFarEnabled=true;App.bGpuFarAuditEnabled=true;App.bShapeWorld=true;App.bForceGenerate=true;App.bQuitAfterClipmap=true;}
         else if(Arg=="--shape-audit")App.bShapeWorld=true,App.bShapeAuditOnly=true;
         else if(Arg=="--grid-stair-audit")bGridStairAudit=true,App.bShapeWorld=true;
         else if(Arg=="--tile-path-hash-audit")bTilePathHashAudit=true,App.bShapeWorld=true;
@@ -4669,7 +4925,7 @@ int main(int Argc,char** Argv)
     GApp=&App;
     if(!App.Window||!App.InitializeDevice()){MessageBoxA(nullptr,"Direct3D 11 initialization failed.","VoxelForge standalone",MB_ICONERROR);return 3;}
     ShowWindow(App.Window,SW_SHOW);UpdateWindow(App.Window);
-    std::printf("Controls: hold right mouse to look; WASD move; Space/Ctrl (E/Q) up/down; Shift faster; wheel changes speed; F wireframe; G toggles legacy/detailed round walls; 1-9 toggle LOD; 0 toggle all levels; N/P change seed; T swaps base/round tunnels; V swaps owner density/shape intervals; Esc quit.\n");
+    std::printf("Controls: hold right mouse to look; WASD move; Space/Ctrl (E/Q) up/down; Shift faster; wheel changes speed; F wireframe; G toggles legacy/detailed round walls; H toggles CPU/GPU far tiles; 1-9 toggle LOD; 0 toggle all levels; N/P change seed; T swaps base/round tunnels; V swaps owner density/shape intervals; Esc quit.\n");
     std::printf("Owner data: %s; starting at first-strate open-space sample (%.2f, %.2f, %.2f) m.\n",
         App.DataPath.c_str(),App.Position.X,App.Position.Y,App.Position.Z);
     std::fflush(stdout);