wip(standalone): GPU step 28 -- skirted near tiles on GPU (correct, but SLOWER); 16 shape workers

Luna: near skirts were CPU-only by implementation gap, not correctness. Ported with
near-resolution refinement and overflow fallback; near-probe preparation split out so
the skirt kernels no longer inline the near field (first attempt: two fxc /O3 compiles
ran 10.5 h and were stopped). Shape compile proven identical at W1/W8/W16.
-SkipCpuDetailBand retired (154/640 skirt audit failures, +174 ms planning).
Audits: hashes W1/W8/W16, walk 5,642/5,642, whole-near 960 tiles + 640 skirt faces
0 sign/topology diffs (one 0.001 cm vertex-key delta), far/thin 0 diffs, both
survivals pass with 2,074 candidates at 7.00 / 7.02 GiB (was ~6.2), ten captures pass.

My runs outside the sandbox (Step28.exe). NOTE the machine was ~25% slower than last
night for every build (step 27 CPU full machine 0.97-1.02 s at night, 1.22-1.29 s now):
  0xFF00: CPU 1.337 / 1.291 / 1.295 s; GPU miss 1.337 s;
          GPU cached 1.663 / 1.424 / 1.617 s (GPU tiles 1,427 / 2,072 / 1,222)
  0xFFFF: CPU 1.221 / 1.231 / 1.236 s (step 27 same session 1.285 / 1.222 s)
          GPU cached 1.582 / 1.550 / 1.416, 1.635 / 1.441 / 1.470 / 1.445 s
          (GPU tiles vary 593-2,074; timestamps up to ~470 ms)
=> GPU route now slower than the CPU route measured alongside it. Step 27 GPU control
   was not clean (both exes load the same precompiled .cso files).

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

Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h

@@ -146,12 +146,9 @@ namespace VoxelForgeCore
         // World-wide walkability is a validation pass, not an input to the compiled field.
         // Keep it enabled for benchmarks/audits; the interactive viewer can defer it.
         bool validatePassageRoutes = true;
-        // Near GPU visual tiles evaluate owner roughness in the compute shader. Their exact
+        // Optional tile sets support attribution of roughness work to GPU and CPU consumers.
-        // sample lattice may omit CPU proxy-band work; CPU density queries still use the exact
-        // analytic field for those omitted samples. CPU visual tiles remain fully packed.
         std::vector<DetailBandTileRegion> gpuVisualDetailTiles;
         std::vector<DetailBandTileRegion> cpuVisualDetailTiles;
-        bool skipCpuDetailBandForGpuVisuals = false;
     };
 
     struct ShapeBuildStats

Standalone/Core/src/ShapeWorldInternal.h

@@ -66,7 +66,6 @@ public:
         uint64 LocalBandSkips=0;
         uint64 GpuVisualCandidates=0,CpuVisualCandidates=0,GpuOnlyCandidates=0;
         uint64 GpuVisualDetailEvaluations=0,GpuOnlyDetailEvaluations=0;
-        uint64 GpuOnlySkippedCandidateMasks=0;
         uint64 DistanceSamples=0,DistanceSampleNanoseconds=0;
         uint64 DetailSamples=0,DetailSampleNanoseconds=0;
         uint64 ReliefSampleNanoseconds=0,LayerLookupSampleNanoseconds=0;
@@ -282,18 +281,14 @@ public:
                bool bInRoundTunnels=false,bool bInProfileDetailBand=false,
                bool bValidatePassageRoutes=true,
                const std::vector<VoxelForgeCore::DetailBandTileRegion>& InGpuVisualDetailTiles={},
-               const std::vector<VoxelForgeCore::DetailBandTileRegion>& InCpuVisualDetailTiles={},
+               const std::vector<VoxelForgeCore::DetailBandTileRegion>& InCpuVisualDetailTiles={})
-               bool bInSkipCpuDetailBandForGpuVisuals=false)
     {
         using Clock=std::chrono::steady_clock;
         const auto TotalStart=Clock::now();
-        Owner=&World;WorkerCount=std::clamp(Workers,1,8);DetailVariant=Variant;
+        Owner=&World;WorkerCount=std::clamp(Workers,1,16);DetailVariant=Variant;
         Scheduler=InScheduler;RoundTunnelsEnabled=bInRoundTunnels;
         bAnalyticOnly=false;
         bProfileDetailBand=bInProfileDetailBand;
-        bSkipCpuDetailBandForGpuVisuals=bInSkipCpuDetailBandForGpuVisuals;
-        if(bSkipCpuDetailBandForGpuVisuals&&InGpuVisualDetailTiles.empty())
-        {Error="GPU visual detail-band skipping requires at least one GPU near tile region";return false;}
         for(int32 A=0;A<3;++A)
         {
             Min[A]=(&MinVoxels.X)[A]-2;
@@ -333,9 +328,6 @@ public:
             Stats.CandidateMergeSeconds);
         if(!BuildGpuVisualDetailNodeSets(InGpuVisualDetailTiles,InCpuVisualDetailTiles,Error))
             return false;
-        GpuVisualProxyValidWords.clear();
-        if(bSkipCpuDetailBandForGpuVisuals)
-            GpuVisualProxyValidWords.assign(CandidateWords.size(),0u);
         Stats.IntervalSeconds=std::chrono::duration<double>(Clock::now()-Stage).count();
 
         Stage=Clock::now();
@@ -777,14 +769,6 @@ public:
                 const uint64 Mask=uint64{1}<<(Zi&63u);
                 if((CandidateWords[Word]&Mask)==0)
                     Out.nearPackedProxy[SampleIndex]=0u;
-                else if(IsSkippedGpuVisualDetailNode(WX,WY,WZ)
-                    &&(GpuVisualProxyValidWords[Word]&Mask)==0)
-                {
-                    // UINT_MAX asks ApplyNearFieldDetailCS to use its own exact
-                    // GPU roughness result instead of a CPU packed proxy.
-                    Out.nearPackedProxy[SampleIndex]=std::numeric_limits<uint32>::max();
-                    ++Out.nearPackedSkippedSampleCount;
-                }
                 else
                 {
                     ++Out.nearPackedCandidateSampleCount;
@@ -933,8 +917,6 @@ public:
     {
         int32 IX=0,IY=0,IZ=0;
         const bool bInteger=NearInteger(X,IX)&&NearInteger(Y,IY)&&NearInteger(Z,IZ);
-        if(bInteger&&IsSkippedGpuVisualDetailNode(IX,IY,IZ))
-            return GpuSkippedDetailDensityAt(IX,IY,IZ,X,Y,Z);
         const bool bPacked=!bAnalyticOnly&&bInteger&&Contains(IX,IY,IZ);
         if(bPacked)
             return PackedDensityAt(IX,IY,IZ,IsRock(IX,IY,IZ),X,Y,Z);
@@ -943,9 +925,6 @@ public:
 
     float DensityAtLattice(int32 X,int32 Y,int32 Z) const
     {
-        if(IsSkippedGpuVisualDetailNode(X,Y,Z))
-            return GpuSkippedDetailDensityAt(X,Y,Z,static_cast<float>(X),
-                static_cast<float>(Y),static_cast<float>(Z));
         const bool bPacked=!bAnalyticOnly&&Contains(X,Y,Z);
         if(bPacked)
             return PackedDensityAt(X,Y,Z,IsRock(X,Y,Z),static_cast<float>(X),
@@ -959,39 +938,6 @@ public:
         return AnalyticDensityAt(X,Y,Z,false);
     }
 
-    float GpuSkippedDetailDensityAt(int32 X,int32 Y,int32 Z,
-                                    float FX,float FY,float FZ) const
-    {
-        const size_t Column=static_cast<size_t>(Y-Min[1])*Size[0]
-            +static_cast<size_t>(X-Min[0]);
-        const size_t Zi=static_cast<size_t>(Z-Min[2]);
-        const size_t Word=Column*WordsPerZ+(Zi>>6);
-        const uint64 Mask=uint64{1}<<(Zi&63u);
-        if(!GpuVisualProxyValidWords.empty()&&(GpuVisualProxyValidWords[Word]&Mask))
-            return PackedDensityAt(X,Y,Z,IsRock(X,Y,Z),FX,FY,FZ);
-
-        if(CandidateWords[Word]&Mask)
-        {
-            // The GPU-only compile omitted this packed cache entry. Reconstruct the
-            // CPU reference's un-clamped detail density and the same signed Q8.8
-            // proxy when an exceptional CPU mesh or field query needs the sample.
-            const float ExactDensity=AnalyticDensityAt(FX,FY,FZ,false,
-                nullptr,nullptr,nullptr,false);
-            int32 Q=static_cast<int32>(std::lround(ExactDensity*256.0f));
-            Q=std::clamp(Q,-32767,32767);
-            if(Q==0)Q=ExactDensity<0.0f?-1:1;
-            float Density=static_cast<float>(Q)*(1.0f/256.0f);
-            const bool bRock=ExactDensity<=0.0f;
-            if((Density<0.0f)!=bRock)
-                Density=(bRock?-1.0f:1.0f)*std::max(std::abs(Density),1.0f/256.0f);
-            return ApplyEditBrushes(Density,FX,FY,FZ);
-        }
-
-        const bool bRock=((RockWords[Word]&Mask)!=0);
-        constexpr float ProxyBandEdge=DetailSupremum+0.125f;
-        return ApplyEditBrushes(bRock?-ProxyBandEdge:ProxyBandEdge,FX,FY,FZ);
-    }
-
     VoxelForgeCore::DetailFieldSample DetailSampleAt(float X,float Y,float Z) const
     {
         VoxelForgeCore::DetailFieldSample Sample;
@@ -1018,18 +964,6 @@ public:
                 OutDensity[Lane]=DensityAtLattice(XStart+Lane*XStep,Y,Z);
             return;
         }
-        if(bSkipCpuDetailBandForGpuVisuals)
-        {
-            bool bContainsSkippedNode=false;
-            for(int32 Lane=0;Lane<Count;++Lane)
-                bContainsSkippedNode|=IsSkippedGpuVisualDetailNode(XStart+Lane*XStep,Y,Z);
-            if(bContainsSkippedNode)
-            {
-                for(int32 Lane=0;Lane<Count;++Lane)
-                    OutDensity[Lane]=DensityAtLattice(XStart+Lane*XStep,Y,Z);
-                return;
-            }
-        }
         bool bFast=EditBrushes.empty()&&Y>=Min[1]&&Y<=Max[1]
             &&Z>=Min[2]&&Z<=Max[2]
             &&XStart>=Min[0]
@@ -1185,9 +1119,6 @@ public:
 
     bool IsRock(int32 X,int32 Y,int32 Z) const
     {
-        if(IsSkippedGpuVisualDetailNode(X,Y,Z))
-            return AnalyticDensityAt(static_cast<float>(X),static_cast<float>(Y),
-                static_cast<float>(Z))<=0.0f;
         if(!Contains(X,Y,Z))return EvaluateFast(X,Y,Z);
         const size_t Line=static_cast<size_t>(Y-Min[1])*Size[0]
             +static_cast<size_t>(X-Min[0]);
@@ -3342,17 +3273,14 @@ private:
         if(GpuTiles.empty()&&CpuTiles.empty())
         {
             GpuVisualDetailWords.clear();CpuVisualDetailWords.clear();
-            CpuVisualDetailGuardWords.clear();
             if(bProfileDetailBand)
             {
                 GpuVisualDetailWords.assign(CandidateWords.size(),0u);
                 CpuVisualDetailWords.assign(CandidateWords.size(),0u);
-                CpuVisualDetailGuardWords.assign(CandidateWords.size(),0u);
             }
             GpuVisualDetailTileCount=CpuVisualDetailTileCount=0;
             GpuVisualDetailGridSamples=CpuVisualDetailGridSamples=0;
             GpuVisualDetailNodeCount=CpuVisualDetailNodeCount=0;
-            CpuVisualDetailGuardNodeCount=0;
             return true;
         }
         const auto BuildSet=[&](const std::vector<VoxelForgeCore::DetailBandTileRegion>& Tiles,
@@ -3403,55 +3331,16 @@ private:
         GpuVisualDetailNodeCount=0;CpuVisualDetailNodeCount=0;
         for(const uint64 Bits:GpuVisualDetailWords)GpuVisualDetailNodeCount+=std::popcount(Bits);
         for(const uint64 Bits:CpuVisualDetailWords)CpuVisualDetailNodeCount+=std::popcount(Bits);
-        CpuVisualDetailGuardWords.assign(CandidateWords.size(),0u);
-        if(bSkipCpuDetailBandForGpuVisuals&&!CpuTiles.empty())
-        {
-            constexpr int32 AxisOffsets[6][3]={{-1,0,0},{1,0,0},{0,-1,0},
-                {0,1,0},{0,0,-1},{0,0,1}};
-            for(size_t Word=0;Word<CpuVisualDetailWords.size();++Word)
-            {
-                uint64 Bits=CpuVisualDetailWords[Word];
-                const size_t Column=Word/WordsPerZ;
-                const int32 X=Min[0]+static_cast<int32>(Column%Size[0]);
-                const int32 Y=Min[1]+static_cast<int32>(Column/Size[0]);
-                const size_t ZiBase=(Word%WordsPerZ)*64u;
-                while(Bits)
-                {
-                    const uint32 Bit=std::countr_zero(Bits);Bits&=Bits-1;
-                    const size_t Zi=ZiBase+Bit;
-                    if(Zi>=static_cast<size_t>(Size[2]))continue;
-                    const int32 Z=Min[2]+static_cast<int32>(Zi);
-                    for(const auto& Offset:AxisOffsets)
-                    {
-                        const int32 GX=X+Offset[0],GY=Y+Offset[1],GZ=Z+Offset[2];
-                        if(GX<Min[0]||GX>Max[0]||GY<Min[1]||GY>Max[1]
-                            ||GZ<Min[2]||GZ>Max[2])continue;
-                        const size_t GuardColumn=static_cast<size_t>(GY-Min[1])*Size[0]
-                            +static_cast<size_t>(GX-Min[0]);
-                        const size_t GuardZi=static_cast<size_t>(GZ-Min[2]);
-                        const size_t GuardWord=GuardColumn*WordsPerZ+(GuardZi>>6);
-                        const uint64 GuardMask=uint64{1}<<(GuardZi&63u);
-                        if(CandidateWords[GuardWord]&GuardMask)
-                            CpuVisualDetailGuardWords[GuardWord]|=GuardMask;
-                    }
-                }
-            }
-        }
-        CpuVisualDetailGuardNodeCount=0;
-        for(const uint64 Bits:CpuVisualDetailGuardWords)
-            CpuVisualDetailGuardNodeCount+=std::popcount(Bits);
         GpuVisualDetailTileCount=GpuTiles.size();CpuVisualDetailTileCount=CpuTiles.size();
         if(!GpuTiles.empty()||!CpuTiles.empty())
         {
-            std::printf("[GpuVisualDetailBandPlan] gpu_tiles=%llu gpu_sample_nodes=%llu gpu_unique_candidate_nodes=%llu cpu_near_tiles=%llu cpu_sample_nodes=%llu cpu_unique_candidate_nodes=%llu cpu_boundary_guard_candidate_nodes=%llu skip_enabled=%d\n",
+            std::printf("[GpuVisualDetailBandPlan] gpu_tiles=%llu gpu_sample_nodes=%llu gpu_unique_candidate_nodes=%llu cpu_near_tiles=%llu cpu_sample_nodes=%llu cpu_unique_candidate_nodes=%llu\n",
                 static_cast<unsigned long long>(GpuVisualDetailTileCount),
                 static_cast<unsigned long long>(GpuVisualDetailGridSamples),
                 static_cast<unsigned long long>(GpuVisualDetailNodeCount),
                 static_cast<unsigned long long>(CpuVisualDetailTileCount),
                 static_cast<unsigned long long>(CpuVisualDetailGridSamples),
-                static_cast<unsigned long long>(CpuVisualDetailNodeCount),
+                static_cast<unsigned long long>(CpuVisualDetailNodeCount));
-                static_cast<unsigned long long>(CpuVisualDetailGuardNodeCount),
-                bSkipCpuDetailBandForGpuVisuals?1:0);
         }
         return true;
     }
@@ -3476,25 +3365,6 @@ private:
             &(uint64{1}<<(Zi&63u)))!=0;
     }
 
-    bool IsSkippedGpuVisualDetailWord(size_t Word,uint64 Mask) const
-    {
-        return bSkipCpuDetailBandForGpuVisuals
-            &&(GpuVisualDetailWords[Word]&Mask)!=0
-            &&(CpuVisualDetailWords[Word]&Mask)==0
-            &&(CpuVisualDetailGuardWords[Word]&Mask)==0;
-    }
-
-    bool IsSkippedGpuVisualDetailNode(int32 X,int32 Y,int32 Z) const
-    {
-        if(!bSkipCpuDetailBandForGpuVisuals||!Contains(X,Y,Z)
-            ||GpuVisualDetailWords.empty())return false;
-        const size_t Column=static_cast<size_t>(Y-Min[1])*Size[0]
-            +static_cast<size_t>(X-Min[0]);
-        const size_t Zi=static_cast<size_t>(Z-Min[2]);
-        return IsSkippedGpuVisualDetailWord(Column*WordsPerZ+(Zi>>6),
-            uint64{1}<<(Zi&63u));
-    }
-
     void ApplyDetailBand(FParallelTiming& Timing)
     {
         using Clock=std::chrono::steady_clock;
@@ -3566,25 +3436,16 @@ private:
                         if(std::abs(SDF)>DetailSupremum+0.125f)continue;
                         const size_t VisualWordIndex=WordBase+(Zi>>6);
                         const uint64 VisualMask=uint64{1}<<(Zi&63u);
-                        const bool bGpuVisualNode=(bProfile||bSkipCpuDetailBandForGpuVisuals)
+                        const bool bGpuVisualNode=bProfile
                             &&(GpuVisualDetailWords[VisualWordIndex]&VisualMask)!=0;
-                        const bool bCpuVisualNode=(bProfile||bSkipCpuDetailBandForGpuVisuals)
+                        const bool bCpuVisualNode=bProfile
                             &&(CpuVisualDetailWords[VisualWordIndex]&VisualMask)!=0;
-                        const bool bSkipDetail=bSkipCpuDetailBandForGpuVisuals&&bGpuVisualNode
-                            &&!bCpuVisualNode
-                            &&(CpuVisualDetailGuardWords[VisualWordIndex]&VisualMask)==0;
                         if(bProfile)
                         {
                             if(bGpuVisualNode)++Profile.GpuVisualCandidates;
                             if(bCpuVisualNode)++Profile.CpuVisualCandidates;
                             if(bGpuVisualNode&&!bCpuVisualNode)++Profile.GpuOnlyCandidates;
                         }
-                        if(bSkipDetail)
-                        {
-                            ValidBits|=uint64{1}<<(Zi&63u);
-                            if(bProfile)++Profile.GpuOnlySkippedCandidateMasks;
-                            continue;
-                        }
                         const float SurfaceSDF=Sample.SurfaceSDF;
                         const float CaveInfluence=Sample.CaveInfluence;
                         const float FloorInfluence=Sample.FloorInfluence;
@@ -3681,7 +3542,7 @@ private:
             uint64 LayerLookupSampleNs=0,OwnerSampleNs=0,OwnerFieldCalls=0,OwnerActiveCalls=0;
             uint64 OwnerCacheHits=0,OwnerCacheMisses=0;
             uint64 GpuVisualCandidates=0,CpuVisualCandidates=0,GpuOnlyCandidates=0;
-            uint64 GpuVisualDetailCalls=0,GpuOnlyDetailCalls=0,GpuOnlySkippedMasks=0;
+            uint64 GpuVisualDetailCalls=0,GpuOnlyDetailCalls=0;
             for(int32 Worker=0;Worker<WorkerCount;++Worker)
             {
                 const FDetailWorkerProfile& P=WorkerProfile[static_cast<size_t>(Worker)];
@@ -3701,7 +3562,6 @@ private:
                 GpuOnlyCandidates+=P.GpuOnlyCandidates;
                 GpuVisualDetailCalls+=P.GpuVisualDetailEvaluations;
                 GpuOnlyDetailCalls+=P.GpuOnlyDetailEvaluations;
-                GpuOnlySkippedMasks+=P.GpuOnlySkippedCandidateMasks;
             }
             const double DistanceAverageNs=DistanceSamples?static_cast<double>(DistanceSampleNs)/DistanceSamples:0.0;
             const double DetailAverageNs=DetailSamples?static_cast<double>(DetailSampleNs)/DetailSamples:0.0;
@@ -3726,14 +3586,12 @@ private:
                 static_cast<unsigned long long>(OwnerCacheHits),
                 static_cast<unsigned long long>(OwnerCacheMisses),
                 static_cast<unsigned long long>(LocalBandSkips));
-            std::printf("[GpuDetailConsumerProfile] gpu_near_proxy_nodes=%llu cpu_near_mesh_nodes=%llu gpu_only_nodes=%llu gpu_near_detail_noise_evaluations=%llu gpu_only_detail_noise_evaluations=%llu gpu_only_skipped_candidate_masks=%llu skip_enabled=%d\n",
+            std::printf("[GpuDetailConsumerProfile] gpu_near_proxy_nodes=%llu cpu_near_mesh_nodes=%llu gpu_only_nodes=%llu gpu_near_detail_noise_evaluations=%llu gpu_only_detail_noise_evaluations=%llu\n",
                 static_cast<unsigned long long>(GpuVisualCandidates),
                 static_cast<unsigned long long>(CpuVisualCandidates),
                 static_cast<unsigned long long>(GpuOnlyCandidates),
                 static_cast<unsigned long long>(GpuVisualDetailCalls),
-                static_cast<unsigned long long>(GpuOnlyDetailCalls),
+                static_cast<unsigned long long>(GpuOnlyDetailCalls));
-                static_cast<unsigned long long>(GpuOnlySkippedMasks),
-                bSkipCpuDetailBandForGpuVisuals?1:0);
             for(int32 Worker=0;Worker<WorkerCount;++Worker)
             {
                 const FDetailWorkerProfile& P=WorkerProfile[static_cast<size_t>(Worker)];
@@ -3769,8 +3627,6 @@ private:
                     +static_cast<size_t>(X-Min[0]);
                 StoreSurfaceDistance(Node,Density);
                 CandidateWords[Word]|=Mask;
-                if(IsSkippedGpuVisualDetailNode(X,Y,Z))
-                    GpuVisualProxyValidWords[Word]|=Mask;
                 ++AddedThisPass;
 
                 const bool bExactRock=Density<0.0f;
@@ -4821,14 +4677,10 @@ private:
     std::unique_ptr<int16[]> SurfaceDistanceQ;
     std::array<std::vector<uint64>,3> CandidateAxisWords;
     std::vector<uint64> CandidateWords;
-    bool bSkipCpuDetailBandForGpuVisuals=false;
     std::vector<uint64> GpuVisualDetailWords,CpuVisualDetailWords;
-    std::vector<uint64> CpuVisualDetailGuardWords;
     uint64 GpuVisualDetailTileCount=0,CpuVisualDetailTileCount=0;
     uint64 GpuVisualDetailGridSamples=0,CpuVisualDetailGridSamples=0;
     uint64 GpuVisualDetailNodeCount=0,CpuVisualDetailNodeCount=0;
-    uint64 CpuVisualDetailGuardNodeCount=0;
-    std::vector<uint64> GpuVisualProxyValidWords;
     FBuildStats Stats;
     bool bAnalyticOnly=false;
 };

Standalone/Core/src/VoxelForgeCore.cpp

@@ -307,7 +307,7 @@ namespace VoxelForgeCore
             static_cast<int32>(std::clamp<std::uint32_t>(request.workerLimit,1,8)),ToNative(request.detail),
             request.includeFarWorld,error,&scheduler,request.roundTunnels,request.profileDetailBand,
             request.validatePassageRoutes,request.gpuVisualDetailTiles,
-            request.cpuVisualDetailTiles,request.skipCpuDetailBandForGpuVisuals);
+            request.cpuVisualDetailTiles);
         if(ok)
         {
             impl_->shape=std::move(shape);

Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldPort.ps1

@@ -1,10 +1,19 @@
 [CmdletBinding()]
-param([switch]$SkipCore,[string]$OutputName='VoxelForgeGpuFarAgreement.exe')
+param([switch]$SkipCore,[string]$OutputName='VoxelForgeGpuFarAgreement.exe',
+    [string]$AffinityMask='0xFF00',[switch]$NormalPriority)
 
 $ErrorActionPreference = 'Stop'
 $HostProcess = Get-Process -Id $PID
-$HostProcess.PriorityClass = 'BelowNormal'
+$AffinityText=$AffinityMask.Trim()
-$HostProcess.ProcessorAffinity = [IntPtr]0xFF00
+if($AffinityText -match '^0[xX]([0-9a-fA-F]+)$'){$AffinityValue=[Convert]::ToUInt64($Matches[1],16)}
+elseif($AffinityText -match '^\d+$'){$AffinityValue=[Convert]::ToUInt64($AffinityText,10)}
+else{throw 'AffinityMask must be a non-zero decimal or hexadecimal processor mask.'}
+$AvailableAffinity=[uint64]$HostProcess.ProcessorAffinity.ToInt64()
+if($AffinityValue-eq0-or$AffinityValue-gt[uint64][long]::MaxValue-or(($AffinityValue-band$AvailableAffinity)-ne$AffinityValue))
+{throw 'AffinityMask must fit the available logical processors.'}
+$BuildPriority=if($NormalPriority){'Normal'}else{'BelowNormal'}
+$HostProcess.PriorityClass=$BuildPriority
+$HostProcess.ProcessorAffinity=[IntPtr][long]$AffinityValue
 $Standalone = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
 $BuildDir = Join-Path $Standalone 'Build\VoxelForgeGpuFarAgreement'
 $WrapperDir = Join-Path $BuildDir 'wrappers'

Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldViewer.ps1

@@ -1,13 +1,24 @@
 [CmdletBinding()]
 param(
     [string]$OutputName = 'VoxelForgeGpuFarHybridViewer.exe',
-    [ValidateSet('All','Core')][string]$Target = 'All'
+    [ValidateSet('All','Core')][string]$Target = 'All',
+    [string]$AffinityMask = '0xFF00',
+    [switch]$NormalPriority,
+    [switch]$SkipShaderCompile
 )
 
 $ErrorActionPreference = 'Stop'
 $BuildHostProcess = Get-Process -Id $PID
-$BuildHostProcess.PriorityClass = 'BelowNormal'
+$AffinityText = $AffinityMask.Trim()
-$BuildHostProcess.ProcessorAffinity = [IntPtr]0xFF00
+if ($AffinityText -match '^0[xX]([0-9a-fA-F]+)$') { $AffinityValue = [Convert]::ToUInt64($Matches[1],16) }
+elseif ($AffinityText -match '^\d+$') { $AffinityValue = [Convert]::ToUInt64($AffinityText,10) }
+else { throw 'AffinityMask must be a non-zero decimal or hexadecimal processor mask.' }
+$AvailableAffinity = [uint64]$BuildHostProcess.ProcessorAffinity.ToInt64()
+if ($AffinityValue -eq 0 -or $AffinityValue -gt [uint64][long]::MaxValue -or (($AffinityValue -band $AvailableAffinity) -ne $AffinityValue))
+{ throw 'AffinityMask must fit the available logical processors.' }
+$BuildPriority = if ($NormalPriority) { 'Normal' } else { 'BelowNormal' }
+$BuildHostProcess.PriorityClass = $BuildPriority
+$BuildHostProcess.ProcessorAffinity = [IntPtr][long]$AffinityValue
 $Root = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
 $BuildStem = [IO.Path]::GetFileNameWithoutExtension($OutputName)
 $BuildDir = Join-Path $Root ('Build\'+$BuildStem)
@@ -113,21 +124,30 @@ if ($Target -eq 'All') {
     $LegacyObjects = @($LegacySources | ForEach-Object { '"{0}.obj"' -f [IO.Path]::GetFileNameWithoutExtension($_) })
     $LegacyRspLines = $CommonFlags + @('/c') + @($LegacySources | ForEach-Object { '"{0}"' -f $_ })
     [IO.File]::WriteAllLines($LegacyRsp,$LegacyRspLines,[Text.UTF8Encoding]::new($false))
-    if (-not (Test-Path -LiteralPath $ShaderSource -PathType Leaf)) { throw "GPU shader source is missing: $ShaderSource" }
+    if ($SkipShaderCompile) {
-    $FxCRoot = 'C:\Program Files (x86)\Windows Kits\10\bin'
+        foreach ($Entry in @('EvaluateFarFieldCS','EvaluateNearFieldCS','ApplyNearFieldDetailCS','ClassifyCS','ClassifyFarCS','ScanGroupsCS','EmitCS','EmitNearCS','NoiseProbeCS','BuildFarConnectorRingsCS','EmitFarConnectorPanelsCS','EvaluateFarSkirtEndpointsCS','PrepareFarSkirtNearProbesCS','FinalizeFarSkirtNearProbeDispatchCS','EvaluateFarSkirtNearProbesCS','PrepareFarSkirtNearChildProbesCS','EvaluateFarSkirtNearChildProbesCS','ClassifyFarSkirtsCS','ScanFarSkirtsCS','EmitFarSkirtsCS')) {
-    $FxC = Get-ChildItem -LiteralPath $FxCRoot -Filter fxc.exe -Recurse -ErrorAction SilentlyContinue |
-        Where-Object { $_.FullName -match '\\x64\\fxc\.exe$' } |
-        Sort-Object -Property FullName -Descending | Select-Object -First 1
-    if ($FxC) {
-        [IO.Directory]::CreateDirectory($PrecompiledShaderDir) | Out-Null
-        foreach ($Entry in @('EvaluateFarFieldCS','EvaluateNearFieldCS','ApplyNearFieldDetailCS','ClassifyCS','ClassifyFarCS','ScanGroupsCS','EmitCS','EmitNearCS','NoiseProbeCS','BuildFarConnectorRingsCS','EmitFarConnectorPanelsCS','ClassifyFarSkirtsCS','ScanFarSkirtsCS','EmitFarSkirtsCS')) {
             $ShaderBinary = Join-Path $PrecompiledShaderDir ($Entry+'.cso')
-            $CommandLine = '"{0}" /nologo /T cs_5_0 /E {1} /O3 /Gis /Ges /Fo "{2}" "{3}"' -f $FxC.FullName,$Entry,$ShaderBinary,$ShaderSource
+            if (-not (Test-Path -LiteralPath $ShaderBinary -PathType Leaf)) {
-            Invoke-FieldViewerBuildCommand $CommandLine "precompile HLSL entry $Entry"
+                throw "Precompiled shader is missing with SkipShaderCompile enabled: $ShaderBinary"
+            }
         }
-        Write-Host "Precompiled viewer compute shaders into $PrecompiledShaderDir"
     } else {
-        Write-Warning 'fxc.exe was not found; the viewer will fall back to one-time runtime shader compilation.'
+        if (-not (Test-Path -LiteralPath $ShaderSource -PathType Leaf)) { throw "GPU shader source is missing: $ShaderSource" }
+        $FxCRoot = 'C:\Program Files (x86)\Windows Kits\10\bin'
+        $FxC = Get-ChildItem -LiteralPath $FxCRoot -Filter fxc.exe -Recurse -ErrorAction SilentlyContinue |
+            Where-Object { $_.FullName -match '\\x64\\fxc\.exe$' } |
+            Sort-Object -Property FullName -Descending | Select-Object -First 1
+        if ($FxC) {
+            [IO.Directory]::CreateDirectory($PrecompiledShaderDir) | Out-Null
+            foreach ($Entry in @('EvaluateFarFieldCS','EvaluateNearFieldCS','ApplyNearFieldDetailCS','ClassifyCS','ClassifyFarCS','ScanGroupsCS','EmitCS','EmitNearCS','NoiseProbeCS','BuildFarConnectorRingsCS','EmitFarConnectorPanelsCS','EvaluateFarSkirtEndpointsCS','PrepareFarSkirtNearProbesCS','FinalizeFarSkirtNearProbeDispatchCS','EvaluateFarSkirtNearProbesCS','PrepareFarSkirtNearChildProbesCS','EvaluateFarSkirtNearChildProbesCS','ClassifyFarSkirtsCS','ScanFarSkirtsCS','EmitFarSkirtsCS')) {
+                $ShaderBinary = Join-Path $PrecompiledShaderDir ($Entry+'.cso')
+                $CommandLine = '"{0}" /nologo /T cs_5_0 /E {1} /O3 /Gis /Ges /Fo "{2}" "{3}"' -f $FxC.FullName,$Entry,$ShaderBinary,$ShaderSource
+                Invoke-FieldViewerBuildCommand $CommandLine "precompile HLSL entry $Entry"
+            }
+            Write-Host "Precompiled viewer compute shaders into $PrecompiledShaderDir"
+        } else {
+            Write-Warning 'fxc.exe was not found; the viewer will fall back to one-time runtime shader compilation.'
+        }
     }
 }
 

Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h

@@ -16,6 +16,8 @@ namespace VoxelForgeGpuFarViewer
     inline constexpr std::uint32_t PooledOutputVertexCapacityPerTile=19456;
     inline constexpr std::uint32_t PooledOutputSlotsPerChunk=128;
     inline constexpr std::uint32_t PooledOutputArgsDwordsPerTile=12;
+    inline constexpr std::uint32_t PooledFarPanelVertexCapacityPerTile=32768;
+    inline constexpr std::uint32_t NearSkirtOutputVertexCapacityPerTile=4096;
     struct CpuVertex { float position[3]{}, normal[3]{}; };
     struct CpuMesh
     {
@@ -43,10 +45,14 @@ namespace VoxelForgeGpuFarViewer
         const std::filesystem::path& shaderPath,std::string& error);
     bool ReserveFarOutputPool(ID3D11Device* device,ID3D11DeviceContext* context,
         const std::filesystem::path& shaderPath,std::uint32_t tileSlots,std::uint32_t cells,
-        std::string& error);
+        std::string& error,std::uint32_t connectorVertexCapacity=UINT32_MAX);
     bool ReadFarTileRequiredVertexCounts(ID3D11Device* device,ID3D11DeviceContext* context,
         const std::filesystem::path& shaderPath,const std::vector<std::uint32_t>& outputSlots,
         std::vector<std::uint32_t>& requiredCounts,std::string& error);
+    bool ReadFarTileOutputStatuses(ID3D11Device* device,ID3D11DeviceContext* context,
+        const std::filesystem::path& shaderPath,const std::vector<std::uint32_t>& outputSlots,
+        std::vector<std::uint32_t>& requiredSurfaceVertexCounts,
+        std::vector<std::uint32_t>& connectorOverflowFlags,std::string& error);
     bool QueryFarGpuMemory(ID3D11Device* device,ID3D11DeviceContext* context,
         const std::filesystem::path& shaderPath,std::uint64_t& usageBytes,
         std::uint64_t& budgetBytes,std::string& error);
@@ -66,7 +72,9 @@ namespace VoxelForgeGpuFarViewer
         ID3D11Query** completionEvent,GpuTimingQueries* timingQueries,
         CpuMesh* auditReadback,
         double* enqueueWallMs, double* gpuDispatchMs, std::string& error,
-        bool captureGpuTiming=true);
+        bool captureGpuTiming=true,
+        std::uint32_t connectorOutputBaseVertex=UINT32_MAX,
+        std::uint32_t connectorOutputCapacity=PooledFarPanelVertexCapacityPerTile);
     bool SubmitFarTileBatch(ID3D11Device* device,ID3D11DeviceContext* immediateContext,
         const std::filesystem::path& shaderPath,ID3D11Query** completionEvent,
         GpuTimingQueries* timingQueries,std::string& error);

Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp

@@ -114,7 +114,7 @@ namespace
         std::uint32_t connectorCount=0,ringCount=0,vertexCount=0,panelVertexCapacity=0;
         std::uint32_t shapeCount=0,skirtFaceMask=0;
         float skirtDepthVoxels=0.0f;
-        std::uint32_t padding=0;
+        std::uint32_t outputBaseVertex=0;
     };
     static_assert(sizeof(FarConnectorConstants)==32);
 
@@ -161,6 +161,15 @@ namespace
         std::uint32_t metadata[4];
     };
     static_assert(sizeof(FarSkirtCachedSamplesGpu) == 40);
+    struct FarSkirtNearProbeGpu
+    {
+        float positionDensity[4];
+    };
+    static_assert(sizeof(FarSkirtNearProbeGpu)==16);
+    static constexpr std::uint32_t FarSkirtNearRootLimit=6;
+    static constexpr std::uint32_t FarSkirtNearProbeStride=43;
+    static constexpr std::size_t FarSkirtNearProbeCapacityPerSurfaceVertex=
+        FarSkirtNearRootLimit*FarSkirtNearProbeStride;
     static_assert(sizeof(FarFieldPrimitiveGpu) == 224);
     static_assert(sizeof(FarFieldSpatialBin) == 24);
     static_assert(sizeof(FarFieldStratumBoundary) == 8);
@@ -530,8 +539,6 @@ namespace
             ComPtr<ID3D11Buffer> argsReadback;
             ComPtr<ID3D11ShaderResourceView> surfaceVertexSrv;
             ComPtr<ID3D11UnorderedAccessView> vertexUav, indexUav, argsUav;
-            ComPtr<ID3D11Buffer> connectorVertices,connectorIndices;
-            ComPtr<ID3D11UnorderedAccessView> connectorVertexUav,connectorIndexUav;
         };
         struct DynamicInput
         {
@@ -551,12 +558,15 @@ namespace
             ComPtr<ID3D11ShaderResourceView> connectorRingSrv;
             ComPtr<ID3D11UnorderedAccessView> densityUav,nearNormalDensityUav,nearDetailBaseUav,caseUav,localOffsetUav,groupCountUav,groupOffsetUav;
             ComPtr<ID3D11UnorderedAccessView> connectorRingUav;
-            ComPtr<ID3D11Buffer> skirtScratch;
+        ComPtr<ID3D11Buffer> skirtScratch;
-            ComPtr<ID3D11UnorderedAccessView> skirtScratchUav;
+        ComPtr<ID3D11UnorderedAccessView> skirtScratchUav;
-            ComPtr<ID3D11Buffer> skirtProbeData;
+        ComPtr<ID3D11Buffer> skirtProbeData,skirtPreparedProbeData,skirtNearProbeData;
-            ComPtr<ID3D11ShaderResourceView> skirtProbeSrv;
+        ComPtr<ID3D11Buffer> skirtNearProbeCounter,skirtNearProbeDispatchArgs;
-            ComPtr<ID3D11UnorderedAccessView> skirtProbeUav;
+        ComPtr<ID3D11ShaderResourceView> skirtProbeSrv,skirtPreparedProbeSrv,skirtNearProbeSrv;
-            std::uint32_t skirtScratchCapacity=0;
+        ComPtr<ID3D11ShaderResourceView> skirtNearProbeCounterSrv;
+        ComPtr<ID3D11UnorderedAccessView> skirtProbeUav,skirtPreparedProbeUav,skirtNearProbeUav;
+        ComPtr<ID3D11UnorderedAccessView> skirtNearProbeCounterUav,skirtNearProbeDispatchArgsUav;
+        std::uint32_t skirtScratchCapacity=0;
             DynamicInput shapes,ranges,ids,protectedAir,protectedSolid,nearPackedProxy;
             DynamicInput connectors;
             std::uint32_t connectorRingCapacity=0;
@@ -571,12 +581,18 @@ namespace
             ComPtr<ID3D11DeviceContext> immediateContext,deferredContext;
         ComPtr<ID3D11ComputeShader> fieldEval, fieldEvalNear, fieldEvalNearDetail, classify, fieldClassify, scan, emit, emitNear, noiseProbe;
         ComPtr<ID3D11ComputeShader> connectorRings,connectorPanels;
+        ComPtr<ID3D11ComputeShader> evaluateFarSkirtEndpoints,prepareFarSkirtNearProbes;
+        ComPtr<ID3D11ComputeShader> finalizeFarSkirtNearProbeDispatch,evaluateFarSkirtNearProbes;
+        ComPtr<ID3D11ComputeShader> prepareFarSkirtNearChildProbes,evaluateFarSkirtNearChildProbes;
         ComPtr<ID3D11ComputeShader> classifyFarSkirts,scanFarSkirts,emitFarSkirts;
         ComPtr<IDXGIAdapter3> memoryAdapter;
         std::string name;
         UINT64 timestampFrequency = 0;
         std::vector<FarOutputChunk> farOutputChunks;
+        ComPtr<ID3D11Buffer> farConnectorVertices,farConnectorIndices;
+        ComPtr<ID3D11UnorderedAccessView> farConnectorVertexUav,farConnectorIndexUav;
         std::uint32_t farOutputCapacitySlots=0;
+        std::uint32_t farConnectorOutputCapacityVertices=0;
         std::uint32_t farOutputSlotsPerChunk=VoxelForgeGpuFarViewer::PooledOutputSlotsPerChunk,
             farMaxVerticesPerTile=0;
         ComPtr<ID3D11Query> shaderWarmupCompletion;
@@ -585,7 +601,8 @@ namespace
         // Full clip-radius-3 skirt audit peaks at 27,840 panel+skirt vertices per
         // far tile; this leaves 4,928 records of headroom without allocating 49k
         // records for every near and far output slot.
-        static constexpr std::uint32_t FarPanelVertexCapacity=32768;
+        static constexpr std::uint32_t FarPanelVertexCapacity=
+            VoxelForgeGpuFarViewer::PooledFarPanelVertexCapacityPerTile;
 
         void CompileShaders(const std::filesystem::path& shaderPath)
         {
@@ -624,6 +641,12 @@ namespace
             Create("NoiseProbeCS",noiseProbe);
             Create("BuildFarConnectorRingsCS",connectorRings);
             Create("EmitFarConnectorPanelsCS",connectorPanels);
+            Create("EvaluateFarSkirtEndpointsCS",evaluateFarSkirtEndpoints);
+            Create("PrepareFarSkirtNearProbesCS",prepareFarSkirtNearProbes);
+            Create("FinalizeFarSkirtNearProbeDispatchCS",finalizeFarSkirtNearProbeDispatch);
+            Create("EvaluateFarSkirtNearProbesCS",evaluateFarSkirtNearProbes);
+            Create("PrepareFarSkirtNearChildProbesCS",prepareFarSkirtNearChildProbes);
+            Create("EvaluateFarSkirtNearChildProbesCS",evaluateFarSkirtNearChildProbes);
             Create("ClassifyFarSkirtsCS",classifyFarSkirts);
             Create("ScanFarSkirtsCS",scanFarSkirts);
             Create("EmitFarSkirtsCS",emitFarSkirts);
@@ -636,10 +659,13 @@ namespace
         {
             const auto WarmupStart=Clock::now();
             const auto ResourceCreateStart=Clock::now();
-            static constexpr std::array<std::string_view,14> entries={
+            static constexpr std::array<std::string_view,20> entries={
                 "EvaluateFarFieldCS","EvaluateNearFieldCS","ClassifyCS","ClassifyFarCS","ScanGroupsCS","EmitCS",
                 "ApplyNearFieldDetailCS",
                 "EmitNearCS","BuildFarConnectorRingsCS","EmitFarConnectorPanelsCS",
+                "EvaluateFarSkirtEndpointsCS","PrepareFarSkirtNearProbesCS",
+                "FinalizeFarSkirtNearProbeDispatchCS","EvaluateFarSkirtNearProbesCS",
+                "PrepareFarSkirtNearChildProbesCS","EvaluateFarSkirtNearChildProbesCS",
                 "ClassifyFarSkirtsCS","ScanFarSkirtsCS","EmitFarSkirtsCS","NoiseProbeCS"};
             const bool skipEntryFilter=entryFilter.starts_with("skip:");
             const std::string_view filteredEntry=skipEntryFilter
@@ -653,6 +679,7 @@ namespace
             tile.nearFieldMode=(entryFilter.empty()||filteredEntry=="EvaluateNearFieldCS"
                 ||filteredEntry=="ApplyNearFieldDetailCS"||filteredEntry=="EmitNearCS")?1u:0u;
             tile.voxelSizeCm=25.0f;tile.outputVertexCapacity=3;
+            tile.padding[0]=3;
             FarFieldConstants field{};
             field.strataLo=-100.0f;field.strataHi=100.0f;field.strataBlend=12.0f;
             field.sampleCount=sampleCount;field.spatialBinMask=0;field.spatialBinSize=64;
@@ -708,6 +735,14 @@ namespace
             auto skirtScratch=Structured<std::uint32_t>(nullptr,3,D3D11_BIND_UNORDERED_ACCESS);
             auto skirtProbeData=Structured<FarSkirtCachedSamplesGpu>(nullptr,3,
                 D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
+            auto skirtPreparedProbeData=Structured<FarSkirtCachedSamplesGpu>(nullptr,3,
+                D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
+            const std::size_t warmupNearProbeCapacity=3u*FarSkirtNearProbeCapacityPerSurfaceVertex;
+            auto skirtNearProbeData=Structured<FarSkirtNearProbeGpu>(nullptr,warmupNearProbeCapacity,
+                D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
+            auto skirtNearProbeCounter=RawMeshBuffer(sizeof(std::uint32_t),
+                D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS,
+                D3D11_RESOURCE_MISC_BUFFER_ALLOW_RAW_VIEWS);
 
             auto densitySrv=Srv<float>(density.Get(),sampleCount);
             auto nearNormalSrv=Srv<float>(nearNormal.Get(),sampleCount);
@@ -730,6 +765,9 @@ namespace
             auto connectorPanelSrv=Srv<FarFieldConnectorPanel>(panelBuffer.Get(),1);
             auto connectorRingSrv=Srv<FarConnectorRingRecord>(connectorRingBuffer.Get(),1);
             auto skirtProbeSrv=Srv<FarSkirtCachedSamplesGpu>(skirtProbeData.Get(),3);
+            auto skirtPreparedProbeSrv=Srv<FarSkirtCachedSamplesGpu>(skirtPreparedProbeData.Get(),3);
+            auto skirtNearProbeSrv=Srv<FarSkirtNearProbeGpu>(skirtNearProbeData.Get(),warmupNearProbeCapacity);
+            auto skirtNearProbeCounterSrv=RawSrv(skirtNearProbeCounter.Get(),1);
             auto densityUav=Uav(density.Get(),sampleCount);
             auto nearNormalUav=Uav(nearNormal.Get(),sampleCount);
             auto nearDetailBaseUav=Uav(nearDetailBase.Get(),sampleCount);
@@ -741,10 +779,17 @@ namespace
             auto connectorRingUav=Uav(connectorRingBuffer.Get(),1);
             auto skirtScratchUav=Uav(skirtScratch.Get(),3);
             auto skirtProbeUav=Uav(skirtProbeData.Get(),3);
+            auto skirtPreparedProbeUav=Uav(skirtPreparedProbeData.Get(),3);
+            auto skirtNearProbeUav=Uav(skirtNearProbeData.Get(),static_cast<UINT>(warmupNearProbeCapacity));
+            auto skirtNearProbeCounterUav=RawUav(skirtNearProbeCounter.Get(),1);
 
             auto args=RawMeshBuffer(VoxelForgeGpuFarViewer::PooledOutputArgsDwordsPerTile
                 *sizeof(std::uint32_t),D3D11_BIND_UNORDERED_ACCESS,
                 D3D11_RESOURCE_MISC_BUFFER_ALLOW_RAW_VIEWS|D3D11_RESOURCE_MISC_DRAWINDIRECT_ARGS);
+            auto skirtNearProbeDispatchArgs=RawMeshBuffer(3u*sizeof(std::uint32_t),
+                D3D11_BIND_UNORDERED_ACCESS,
+                D3D11_RESOURCE_MISC_BUFFER_ALLOW_RAW_VIEWS|D3D11_RESOURCE_MISC_DRAWINDIRECT_ARGS);
+            auto skirtNearProbeDispatchArgsUav=RawUav(skirtNearProbeDispatchArgs.Get(),3);
             auto vertices=RawMeshBuffer(3u*sizeof(GpuVertex),D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_SHADER_RESOURCE);
             auto indices=RawMeshBuffer(3u*sizeof(std::uint32_t),D3D11_BIND_UNORDERED_ACCESS);
             auto panelVertices=RawMeshBuffer(64u*sizeof(GpuVertex),D3D11_BIND_UNORDERED_ACCESS);
@@ -788,9 +833,9 @@ namespace
             ID3D11Buffer* constants[3]={tileCb.Get(),fieldCb.Get(),connectorCb.Get()};
             context->CSSetConstantBuffers(0,3,constants);
             std::uint32_t dispatchCount=0;
-            auto ClearBindings=[&](UINT srvCount=22)
+            auto ClearBindings=[&](UINT srvCount=24)
             {
-                ID3D11ShaderResourceView* nullSrvs[22]{};
+                ID3D11ShaderResourceView* nullSrvs[24]{};
                 ID3D11UnorderedAccessView* nullUavs[8]{};
                 context->CSSetShaderResources(0,srvCount,nullSrvs);
                 context->CSSetUnorderedAccessViews(0,8,nullUavs,nullptr);
@@ -916,15 +961,71 @@ namespace
             context->CSSetUnorderedAccessViews(0,8,panelUavs,nullptr);
             Dispatch(connectorPanels.Get(),"EmitFarConnectorPanelsCS",stateStart);ClearBindings();
 
-            ID3D11ShaderResourceView* skirtSrvs[22]{};
+            ID3D11ShaderResourceView* skirtSrvs[24]{};
             skirtSrvs[10]=shapeSrv.Get();skirtSrvs[11]=binSrv.Get();skirtSrvs[12]=shapeIdSrv.Get();
             skirtSrvs[13]=protectedAirSrv.Get();skirtSrvs[14]=protectedSolidSrv.Get();skirtSrvs[19]=vertexSrv.Get();
+            skirtSrvs[16]=detailSrv.Get();skirtSrvs[17]=proxySrv.Get();
+            std::array<UINT,4> zeroCounter{};
+            context->ClearUnorderedAccessViewUint(skirtNearProbeCounterUav.Get(),zeroCounter.data());
+            stateStart=Clock::now();
+            context->CSSetShaderResources(0,24,skirtSrvs);
+            ID3D11UnorderedAccessView* endpointSkirtUavs[8]{};
+            endpointSkirtUavs[2]=skirtProbeUav.Get();endpointSkirtUavs[4]=argsUav.Get();
+            context->CSSetUnorderedAccessViews(0,8,endpointSkirtUavs,nullptr);
+            Dispatch(evaluateFarSkirtEndpoints.Get(),"EvaluateFarSkirtEndpointsCS",stateStart);ClearBindings();
+
+            stateStart=Clock::now();
             skirtSrvs[21]=skirtProbeSrv.Get();
+            ID3D11UnorderedAccessView* prepareSkirtUavs[8]{};
+            prepareSkirtUavs[2]=skirtPreparedProbeUav.Get();prepareSkirtUavs[3]=skirtNearProbeUav.Get();
+            prepareSkirtUavs[4]=argsUav.Get();prepareSkirtUavs[5]=skirtNearProbeCounterUav.Get();
+            context->CSSetShaderResources(0,24,skirtSrvs);
+            context->CSSetUnorderedAccessViews(0,8,prepareSkirtUavs,nullptr);
+            Dispatch(prepareFarSkirtNearProbes.Get(),"PrepareFarSkirtNearProbesCS",stateStart);ClearBindings();
+
+            stateStart=Clock::now();
+            ID3D11ShaderResourceView* countSrvs[24]{};countSrvs[23]=skirtNearProbeCounterSrv.Get();
+            ID3D11UnorderedAccessView* dispatchArgsUavs[8]{};
+            dispatchArgsUavs[6]=skirtNearProbeDispatchArgsUav.Get();
+            context->CSSetShaderResources(0,24,countSrvs);
+            context->CSSetUnorderedAccessViews(0,8,dispatchArgsUavs,nullptr);
+            Dispatch(finalizeFarSkirtNearProbeDispatch.Get(),
+                "FinalizeFarSkirtNearProbeDispatchCS",stateStart);ClearBindings();
+
+            stateStart=Clock::now();
+            skirtSrvs[21]=nullptr;skirtSrvs[22]=nullptr;
+            skirtSrvs[23]=skirtNearProbeCounterSrv.Get();
+            ID3D11UnorderedAccessView* evalSkirtUavs[8]{};evalSkirtUavs[3]=skirtNearProbeUav.Get();
+            context->CSSetShaderResources(0,24,skirtSrvs);
+            context->CSSetUnorderedAccessViews(0,8,evalSkirtUavs,nullptr);
+            Dispatch(evaluateFarSkirtNearProbes.Get(),"EvaluateFarSkirtNearProbesCS",stateStart);ClearBindings();
+
+            stateStart=Clock::now();
+            skirtSrvs[21]=skirtPreparedProbeSrv.Get();skirtSrvs[22]=nullptr;skirtSrvs[23]=nullptr;
+            ID3D11UnorderedAccessView* prepareChildUavs[8]{};
+            prepareChildUavs[3]=skirtNearProbeUav.Get();
+            prepareChildUavs[4]=argsUav.Get();
+            context->CSSetShaderResources(0,24,skirtSrvs);
+            context->CSSetUnorderedAccessViews(0,8,prepareChildUavs,nullptr);
+            Dispatch(prepareFarSkirtNearChildProbes.Get(),
+                "PrepareFarSkirtNearChildProbesCS",stateStart);ClearBindings();
+
+            stateStart=Clock::now();
+            ID3D11ShaderResourceView* evalChildSrvs[24]{};
+            evalChildSrvs[23]=skirtNearProbeCounterSrv.Get();
+            ID3D11UnorderedAccessView* evalChildUavs[8]{};evalChildUavs[3]=skirtNearProbeUav.Get();
+            context->CSSetShaderResources(0,24,evalChildSrvs);
+            context->CSSetUnorderedAccessViews(0,8,evalChildUavs,nullptr);
+            Dispatch(evaluateFarSkirtNearChildProbes.Get(),
+                "EvaluateFarSkirtNearChildProbesCS",stateStart);ClearBindings();
+
             stateStart=Clock::now();
-            context->CSSetShaderResources(0,22,skirtSrvs);
             ID3D11UnorderedAccessView* classifySkirtUavs[8]{};
             classifySkirtUavs[0]=skirtScratchUav.Get();classifySkirtUavs[2]=skirtProbeUav.Get();
             classifySkirtUavs[4]=argsUav.Get();
+            skirtSrvs[21]=skirtPreparedProbeSrv.Get();
+            skirtSrvs[22]=skirtNearProbeSrv.Get();skirtSrvs[23]=nullptr;
+            context->CSSetShaderResources(0,24,skirtSrvs);
             context->CSSetUnorderedAccessViews(0,8,classifySkirtUavs,nullptr);
             Dispatch(classifyFarSkirts.Get(),"ClassifyFarSkirtsCS",stateStart);ClearBindings();
 
@@ -935,7 +1036,8 @@ namespace
             Dispatch(scanFarSkirts.Get(),"ScanFarSkirtsCS",stateStart);ClearBindings();
 
             stateStart=Clock::now();
-            context->CSSetShaderResources(0,22,skirtSrvs);
+            skirtSrvs[21]=skirtProbeSrv.Get();
+            context->CSSetShaderResources(0,24,skirtSrvs);
             ID3D11UnorderedAccessView* emitSkirtUavs[8]{};
             emitSkirtUavs[0]=skirtScratchUav.Get();emitSkirtUavs[4]=argsUav.Get();
             emitSkirtUavs[5]=panelVertexUav.Get();emitSkirtUavs[6]=panelIndexUav.Get();
@@ -1176,12 +1278,24 @@ namespace
             return result;
         }
 
-        void EnsureFarOutputPool(std::uint32_t tileSlots,std::uint32_t cells)
+        void EnsureFarOutputPool(std::uint32_t tileSlots,std::uint32_t cells,
+            std::uint32_t connectorVertexCapacity=std::numeric_limits<std::uint32_t>::max())
         {
             if(tileSlots==0||cells==0)throw std::runtime_error("far output pool dimensions must be non-zero");
-            if(tileSlots<=farOutputCapacitySlots&&!farOutputChunks.empty())return;
+            if(connectorVertexCapacity==std::numeric_limits<std::uint32_t>::max())
+            {
+                const std::uint64_t fullCapacity=static_cast<std::uint64_t>(tileSlots)*FarPanelVertexCapacity;
+                if(fullCapacity>std::numeric_limits<std::uint32_t>::max())
+                    throw std::runtime_error("far connector output pool exceeds the 32-bit vertex limit");
+                connectorVertexCapacity=static_cast<std::uint32_t>(fullCapacity);
+            }
+            if(tileSlots<=farOutputCapacitySlots
+                &&connectorVertexCapacity<=farConnectorOutputCapacityVertices
+                &&!farOutputChunks.empty())return;
             farMaxVerticesPerTile=FarOutputVertexCapacityPerTile;
             farOutputCapacitySlots=std::max(tileSlots,farOutputCapacitySlots);
+            farConnectorOutputCapacityVertices=std::max(connectorVertexCapacity,
+                farConnectorOutputCapacityVertices);
             const std::uint32_t chunkCount=(farOutputCapacitySlots+farOutputSlotsPerChunk-1)/farOutputSlotsPerChunk;
             std::vector<FarOutputChunk> chunks;
             chunks.reserve(chunkCount);
@@ -1207,27 +1321,36 @@ namespace
                 chunk.surfaceVertexSrv=RawSrv(chunk.vertices.Get(),static_cast<UINT>(vertexBytes/sizeof(std::uint32_t)));
                 chunk.indexUav=RawUav(chunk.indices.Get(),static_cast<UINT>(indexBytes/sizeof(std::uint32_t)));
                 chunk.argsUav=RawUav(chunk.args.Get(),static_cast<UINT>(argsBytes/sizeof(std::uint32_t)));
-                const std::uint64_t connectorVertexBytes=static_cast<std::uint64_t>(slots)
-                    *FarPanelVertexCapacity*sizeof(GpuVertex);
-                const std::uint64_t connectorIndexBytes=static_cast<std::uint64_t>(slots)
-                    *FarPanelVertexCapacity*sizeof(std::uint32_t);
-                chunk.connectorVertices=RawMeshBuffer(static_cast<UINT>(connectorVertexBytes),
-                    D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_VERTEX_BUFFER);
-                chunk.connectorIndices=RawMeshBuffer(static_cast<UINT>(connectorIndexBytes),
-                    D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_INDEX_BUFFER);
-                chunk.connectorVertexUav=RawUav(chunk.connectorVertices.Get(),
-                    static_cast<UINT>(connectorVertexBytes/sizeof(std::uint32_t)));
-                chunk.connectorIndexUav=RawUav(chunk.connectorIndices.Get(),
-                    static_cast<UINT>(connectorIndexBytes/sizeof(std::uint32_t)));
                 chunks.push_back(std::move(chunk));
             }
+            const std::uint64_t connectorVertexBytes=static_cast<std::uint64_t>(
+                std::max(1u,farConnectorOutputCapacityVertices))*sizeof(GpuVertex);
+            const std::uint64_t connectorIndexBytes=static_cast<std::uint64_t>(
+                std::max(1u,farConnectorOutputCapacityVertices))*sizeof(std::uint32_t);
+            if(connectorVertexBytes>std::numeric_limits<UINT>::max()
+                ||connectorIndexBytes>std::numeric_limits<UINT>::max())
+                throw std::runtime_error("far connector output pool exceeds the D3D11 buffer size limit");
+            auto connectorVertices=RawMeshBuffer(static_cast<UINT>(connectorVertexBytes),
+                D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_VERTEX_BUFFER);
+            auto connectorIndices=RawMeshBuffer(static_cast<UINT>(connectorIndexBytes),
+                D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_INDEX_BUFFER);
+            auto connectorVertexUav=RawUav(connectorVertices.Get(),
+                static_cast<UINT>(connectorVertexBytes/sizeof(std::uint32_t)));
+            auto connectorIndexUav=RawUav(connectorIndices.Get(),
+                static_cast<UINT>(connectorIndexBytes/sizeof(std::uint32_t)));
             farOutputChunks=std::move(chunks);
+            farConnectorVertices=std::move(connectorVertices);
+            farConnectorIndices=std::move(connectorIndices);
+            farConnectorVertexUav=std::move(connectorVertexUav);
+            farConnectorIndexUav=std::move(connectorIndexUav);
         }
 
-        void ReadFarOutputRequiredVertexCounts(const std::vector<std::uint32_t>& outputSlots,
+        void ReadFarOutputStatuses(const std::vector<std::uint32_t>& outputSlots,
-            std::vector<std::uint32_t>& requiredCounts)
+            std::vector<std::uint32_t>& requiredCounts,
+            std::vector<std::uint32_t>& connectorOverflowFlags)
         {
             requiredCounts.assign(outputSlots.size(),0u);
+            connectorOverflowFlags.assign(outputSlots.size(),0u);
             std::vector<std::vector<std::size_t>> indicesByChunk(farOutputChunks.size());
             for(std::size_t i=0;i<outputSlots.size();++i)
             {
@@ -1249,8 +1372,10 @@ namespace
                 for(const std::size_t outputIndex:indices)
                 {
                     const std::uint32_t localSlot=outputSlots[outputIndex]%farOutputSlotsPerChunk;
-                    requiredCounts[outputIndex]=args[
+                    const std::size_t argBase=static_cast<std::size_t>(localSlot)
-                        localSlot*VoxelForgeGpuFarViewer::PooledOutputArgsDwordsPerTile+11u];
+                        *VoxelForgeGpuFarViewer::PooledOutputArgsDwordsPerTile;
+                    connectorOverflowFlags[outputIndex]=args[argBase+10u];
+                    requiredCounts[outputIndex]=args[argBase+11u];
                 }
                 context->Unmap(chunk.argsReadback.Get(),0);
             }
@@ -1335,6 +1460,18 @@ namespace
                 D3D11_BIND_UNORDERED_ACCESS);
             scratch.skirtProbeData=Structured<FarSkirtCachedSamplesGpu>(nullptr,
                 scratch.skirtScratchCapacity,D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
+            scratch.skirtPreparedProbeData=Structured<FarSkirtCachedSamplesGpu>(nullptr,
+                scratch.skirtScratchCapacity,D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
+            const std::size_t nearProbeCapacity=static_cast<std::size_t>(scratch.skirtScratchCapacity)
+                *FarSkirtNearProbeCapacityPerSurfaceVertex;
+            scratch.skirtNearProbeData=Structured<FarSkirtNearProbeGpu>(nullptr,nearProbeCapacity,
+                D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
+            scratch.skirtNearProbeCounter=RawMeshBuffer(sizeof(std::uint32_t),
+                D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS,
+                D3D11_RESOURCE_MISC_BUFFER_ALLOW_RAW_VIEWS);
+            scratch.skirtNearProbeDispatchArgs=RawMeshBuffer(3u*sizeof(std::uint32_t),
+                D3D11_BIND_UNORDERED_ACCESS,
+                D3D11_RESOURCE_MISC_BUFFER_ALLOW_RAW_VIEWS|D3D11_RESOURCE_MISC_DRAWINDIRECT_ARGS);
             scratch.signSrv=Srv<std::uint32_t>(scratch.signs.Get(),1);
             scratch.triSrv=Srv<std::int32_t>(scratch.tri.Get(),256u*16u);
             scratch.densitySrv=Srv<float>(scratch.density.Get(),samples);
@@ -1357,6 +1494,16 @@ namespace
             scratch.skirtProbeSrv=Srv<FarSkirtCachedSamplesGpu>(scratch.skirtProbeData.Get(),
                 scratch.skirtScratchCapacity);
             scratch.skirtProbeUav=Uav(scratch.skirtProbeData.Get(),scratch.skirtScratchCapacity);
+            scratch.skirtPreparedProbeSrv=Srv<FarSkirtCachedSamplesGpu>(
+                scratch.skirtPreparedProbeData.Get(),scratch.skirtScratchCapacity);
+            scratch.skirtPreparedProbeUav=Uav(scratch.skirtPreparedProbeData.Get(),
+                scratch.skirtScratchCapacity);
+            scratch.skirtNearProbeSrv=Srv<FarSkirtNearProbeGpu>(scratch.skirtNearProbeData.Get(),
+                nearProbeCapacity);
+            scratch.skirtNearProbeUav=Uav(scratch.skirtNearProbeData.Get(),nearProbeCapacity);
+            scratch.skirtNearProbeCounterSrv=RawSrv(scratch.skirtNearProbeCounter.Get(),1);
+            scratch.skirtNearProbeCounterUav=RawUav(scratch.skirtNearProbeCounter.Get(),1);
+            scratch.skirtNearProbeDispatchArgsUav=RawUav(scratch.skirtNearProbeDispatchArgs.Get(),3);
             const auto DynamicConstant=[&](UINT bytes,const char* label)
             {
                 D3D11_BUFFER_DESC desc{};desc.ByteWidth=bytes;desc.Usage=D3D11_USAGE_DYNAMIC;
@@ -1436,7 +1583,9 @@ namespace
                       bool waitForGpu = true, bool useIndirectOutputCount = false,
                       std::uint32_t outputSlot = 0, bool usePooledFarOutput = false,
                       bool captureGpuTiming = true,bool batchTimingMode=false,
-                      bool profileGpuKernels=false)
+                      bool profileGpuKernels=false,
+                      std::uint32_t connectorOutputBaseVertex=std::numeric_limits<std::uint32_t>::max(),
+                      std::uint32_t connectorOutputCapacity=FarPanelVertexCapacity)
         {
             using namespace std::chrono;
             const auto wallStart = Clock::now();
@@ -1469,7 +1618,11 @@ namespace
                 outputBaseVertex=localSlot*farMaxVerticesPerTile;
                 outputArgsByteOffset=localSlot*VoxelForgeGpuFarViewer::PooledOutputArgsDwordsPerTile
                     *sizeof(std::uint32_t);
+                if(connectorOutputBaseVertex==std::numeric_limits<std::uint32_t>::max())
+                    connectorOutputBaseVertex=outputSlot*FarPanelVertexCapacity;
             }
+            else if(connectorOutputBaseVertex==std::numeric_limits<std::uint32_t>::max())
+                connectorOutputBaseVertex=0;
             if (farField && (farField->sampleDim != static_cast<std::int32_t>(tile.dim)
                 || farField->cells != static_cast<std::int32_t>(tile.cells)
                 || farField->step != static_cast<std::int32_t>(tile.step)))
@@ -1641,10 +1794,13 @@ namespace
             ComPtr<ID3D11UnorderedAccessView> connectorVertexUav,connectorIndexUav;
             ComPtr<ID3D11Buffer> skirtScratchBuffer;
             ComPtr<ID3D11UnorderedAccessView> skirtScratchUav;
-            ComPtr<ID3D11Buffer> skirtProbeBuffer;
+            ComPtr<ID3D11Buffer> skirtProbeBuffer,skirtPreparedProbeBuffer,skirtNearProbeBuffer;
-            ComPtr<ID3D11ShaderResourceView> skirtProbeSrv;
+            ComPtr<ID3D11Buffer> skirtNearProbeCounterBuffer,skirtNearProbeDispatchArgsBuffer;
-            ComPtr<ID3D11UnorderedAccessView> skirtProbeUav;
+            ComPtr<ID3D11ShaderResourceView> skirtProbeSrv,skirtPreparedProbeSrv,skirtNearProbeSrv;
-            const std::uint32_t connectorPanelCapacity=pooledChunk?FarPanelVertexCapacity:
+            ComPtr<ID3D11ShaderResourceView> skirtNearProbeCounterSrv;
+            ComPtr<ID3D11UnorderedAccessView> skirtProbeUav,skirtPreparedProbeUav,skirtNearProbeUav;
+            ComPtr<ID3D11UnorderedAccessView> skirtNearProbeCounterUav,skirtNearProbeDispatchArgsUav;
+            const std::uint32_t connectorPanelCapacity=pooledChunk?connectorOutputCapacity:
                 std::max<std::uint32_t>(1,farField&&(farField->visibleSkirtFaceMask!=0)
                     ?FarPanelVertexCapacity:(farField?farField->connectorVertexCount:0));
             if(farField&&farField->connectorVertexCount>connectorPanelCapacity)
@@ -1654,11 +1810,22 @@ namespace
                 args=pooledChunk->args;vertices=pooledChunk->vertices;indices=pooledChunk->indices;
                 surfaceVertexSrv=pooledChunk->surfaceVertexSrv;
                 argsUav=pooledChunk->argsUav;vertexUav=pooledChunk->vertexUav;indexUav=pooledChunk->indexUav;
-                connectorVertices=pooledChunk->connectorVertices;connectorIndices=pooledChunk->connectorIndices;
+                connectorVertices=farConnectorVertices;connectorIndices=farConnectorIndices;
-                connectorVertexUav=pooledChunk->connectorVertexUav;connectorIndexUav=pooledChunk->connectorIndexUav;
+                connectorVertexUav=farConnectorVertexUav;connectorIndexUav=farConnectorIndexUav;
                 skirtScratchBuffer=farScratch.skirtScratch;skirtScratchUav=farScratch.skirtScratchUav;
                 skirtProbeBuffer=farScratch.skirtProbeData;skirtProbeSrv=farScratch.skirtProbeSrv;
                 skirtProbeUav=farScratch.skirtProbeUav;
+                skirtPreparedProbeBuffer=farScratch.skirtPreparedProbeData;
+                skirtPreparedProbeSrv=farScratch.skirtPreparedProbeSrv;
+                skirtPreparedProbeUav=farScratch.skirtPreparedProbeUav;
+                skirtNearProbeBuffer=farScratch.skirtNearProbeData;
+                skirtNearProbeSrv=farScratch.skirtNearProbeSrv;
+                skirtNearProbeUav=farScratch.skirtNearProbeUav;
+                skirtNearProbeCounterBuffer=farScratch.skirtNearProbeCounter;
+                skirtNearProbeCounterSrv=farScratch.skirtNearProbeCounterSrv;
+                skirtNearProbeCounterUav=farScratch.skirtNearProbeCounterUav;
+                skirtNearProbeDispatchArgsBuffer=farScratch.skirtNearProbeDispatchArgs;
+                skirtNearProbeDispatchArgsUav=farScratch.skirtNearProbeDispatchArgsUav;
             }
             else
             {
@@ -1690,17 +1857,45 @@ namespace
                     skirtProbeSrv=Srv<FarSkirtCachedSamplesGpu>(skirtProbeBuffer.Get(),
                         static_cast<UINT>(maxVertices));
                     skirtProbeUav=Uav(skirtProbeBuffer.Get(),static_cast<UINT>(maxVertices));
+                    skirtPreparedProbeBuffer=Structured<FarSkirtCachedSamplesGpu>(nullptr,
+                        static_cast<UINT>(maxVertices),
+                        D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
+                    skirtPreparedProbeSrv=Srv<FarSkirtCachedSamplesGpu>(skirtPreparedProbeBuffer.Get(),
+                        static_cast<UINT>(maxVertices));
+                    skirtPreparedProbeUav=Uav(skirtPreparedProbeBuffer.Get(),
+                        static_cast<UINT>(maxVertices));
+                    const std::size_t nearProbeSurfaceCapacity=farField->nearDetailEnabled
+                        &&tile.step<=2?std::max<std::size_t>(1,tile.expectedVertexCount):1u;
+                    const std::size_t nearProbeCapacity=nearProbeSurfaceCapacity
+                        *FarSkirtNearProbeCapacityPerSurfaceVertex;
+                    skirtNearProbeBuffer=Structured<FarSkirtNearProbeGpu>(nullptr,nearProbeCapacity,
+                        D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
+                    skirtNearProbeSrv=Srv<FarSkirtNearProbeGpu>(skirtNearProbeBuffer.Get(),nearProbeCapacity);
+                    skirtNearProbeUav=Uav(skirtNearProbeBuffer.Get(),nearProbeCapacity);
+                    skirtNearProbeCounterBuffer=RawMeshBuffer(sizeof(std::uint32_t),
+                        D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS,
+                        D3D11_RESOURCE_MISC_BUFFER_ALLOW_RAW_VIEWS);
+                    skirtNearProbeCounterSrv=RawSrv(skirtNearProbeCounterBuffer.Get(),1);
+                    skirtNearProbeCounterUav=RawUav(skirtNearProbeCounterBuffer.Get(),1);
+                    skirtNearProbeDispatchArgsBuffer=RawMeshBuffer(3u*sizeof(std::uint32_t),
+                        D3D11_BIND_UNORDERED_ACCESS,
+                        D3D11_RESOURCE_MISC_BUFFER_ALLOW_RAW_VIEWS|D3D11_RESOURCE_MISC_DRAWINDIRECT_ARGS);
+                    skirtNearProbeDispatchArgsUav=RawUav(skirtNearProbeDispatchArgsBuffer.Get(),3);
                 }
             }
             D3D11_BUFFER_DESC cbDesc{};
             cbDesc.ByteWidth = sizeof(TileConstants);
             cbDesc.Usage = D3D11_USAGE_DEFAULT;
             cbDesc.BindFlags = D3D11_BIND_CONSTANT_BUFFER;
+            const std::uint32_t nearProbeSurfaceCapacity=pooledChunk?farMaxVerticesPerTile:
+                (farField&&farField->nearDetailEnabled&&tile.step<=2
+                    ?std::max<std::uint32_t>(1,tile.expectedVertexCount):1u);
             TileConstants constants{tile.cells, tile.dim, cellCount, groupCount,
                 tile.origin[0], tile.origin[1], tile.origin[2], static_cast<std::int32_t>(tile.step),
                 25.0f,outputBaseVertex,outputArgsByteOffset,
                 farField&&farField->nearDetailEnabled?1u:0u,
-                pooledChunk?farMaxVerticesPerTile:static_cast<std::uint32_t>(maxVertices),{0,0,0}};
+                pooledChunk?farMaxVerticesPerTile:static_cast<std::uint32_t>(maxVertices),
+                {nearProbeSurfaceCapacity,0,0}};
             result.outputBaseVertex=outputBaseVertex;
             result.outputArgsByteOffset=outputArgsByteOffset;
             ComPtr<ID3D11Buffer> cb;
@@ -1723,7 +1918,7 @@ namespace
                     farField->connectorRingCount,farField->connectorVertexCount,connectorPanelCapacity,
                     static_cast<std::uint32_t>(farField->worldInput->shapes.size()),
                     static_cast<std::uint32_t>(farField->visibleSkirtFaceMask),
-                    farField->skirtDepthVoxels,0};
+                    farField->skirtDepthVoxels,connectorOutputBaseVertex};
                 if(pooledChunk)
                 {
                     cb=farScratch.tileConstants;farCb=farScratch.farConstants;
@@ -1927,24 +2122,109 @@ namespace
             {
                 ID3D11Buffer* skirtCbs[3]={cb.Get(),farCb.Get(),farConnectorConstantsBuffer.Get()};
                 context->CSSetConstantBuffers(0,3,skirtCbs);
-                ID3D11ShaderResourceView* skirtSrvs[22]{};
+                ID3D11ShaderResourceView* skirtSrvs[24]{};
                 skirtSrvs[10]=farShapeSrv.Get();skirtSrvs[11]=farRangeSrv.Get();
                 skirtSrvs[12]=farIdSrv.Get();skirtSrvs[13]=farProtectedSrv.Get();
-                skirtSrvs[14]=farProtectedSolidSrv.Get();skirtSrvs[19]=surfaceVertexSrv.Get();
+                skirtSrvs[14]=farProtectedSolidSrv.Get();skirtSrvs[17]=farNearProxySrv.Get();
-                skirtSrvs[21]=skirtProbeSrv.Get();
+                skirtSrvs[19]=surfaceVertexSrv.Get();
-                context->CSSetShaderResources(0,22,skirtSrvs);
+                skirtSrvs[16]=farWorldResources.detailLayers.srv.Get();
+                skirtSrvs[23]=skirtNearProbeCounterSrv.Get();
                 ID3D11UnorderedAccessView* classifySkirtUavs[8]{};
                 classifySkirtUavs[0]=skirtScratchUav.Get();classifySkirtUavs[2]=skirtProbeUav.Get();
                 classifySkirtUavs[4]=argsUav.Get();
                 const std::uint32_t skirtDispatchVertices=constants.outputVertexCapacity;
                 beginKernel(6);
+
+                if(constants.nearFieldMode!=0&&constants.stepVoxels<=2)
+                {
+                    const UINT clearWords[4]={0,0,0,0};
+                    context->ClearUnorderedAccessViewUint(skirtNearProbeCounterUav.Get(),clearWords);
+
+                    // Keep the expensive analytic field evaluator in simple one-sample kernels.
+                    // Geometry and topology kernels below consume their cached results.
+                    skirtSrvs[21]=nullptr;skirtSrvs[22]=nullptr;skirtSrvs[23]=nullptr;
+                    ID3D11UnorderedAccessView* endpointUavs[8]{};
+                    endpointUavs[2]=skirtProbeUav.Get();endpointUavs[4]=argsUav.Get();
+                    context->CSSetShaderResources(0,24,skirtSrvs);
+                    context->CSSetShader(evaluateFarSkirtEndpoints.Get(),nullptr,0);
+                    context->CSSetUnorderedAccessViews(0,8,endpointUavs,nullptr);
+                    context->Dispatch((skirtDispatchVertices+63u)/64u,1,1);
+                    ID3D11ShaderResourceView* clearSkirtSrvs[24]{};
+                    ID3D11UnorderedAccessView* clearSkirtUavs[8]{};
+                    context->CSSetShaderResources(0,24,clearSkirtSrvs);
+                    context->CSSetUnorderedAccessViews(0,8,clearSkirtUavs,nullptr);
+
+                    skirtSrvs[21]=skirtProbeSrv.Get();
+                    skirtSrvs[22]=nullptr;skirtSrvs[23]=nullptr;
+                    ID3D11UnorderedAccessView* prepareUavs[8]{};
+                    prepareUavs[2]=skirtPreparedProbeUav.Get();
+                    prepareUavs[3]=skirtNearProbeUav.Get();
+                    prepareUavs[4]=argsUav.Get();
+                    prepareUavs[5]=skirtNearProbeCounterUav.Get();
+                    context->CSSetShaderResources(0,24,skirtSrvs);
+                    context->CSSetShader(prepareFarSkirtNearProbes.Get(),nullptr,0);
+                    context->CSSetUnorderedAccessViews(0,8,prepareUavs,nullptr);
+                    context->Dispatch((skirtDispatchVertices+63u)/64u,1,1);
+                    context->CSSetShaderResources(0,24,clearSkirtSrvs);
+                    context->CSSetUnorderedAccessViews(0,8,clearSkirtUavs,nullptr);
+
+                    ID3D11ShaderResourceView* countSrvs[24]{};
+                    countSrvs[23]=skirtNearProbeCounterSrv.Get();
+                    ID3D11UnorderedAccessView* dispatchArgsUavs[8]{};
+                    dispatchArgsUavs[6]=skirtNearProbeDispatchArgsUav.Get();
+                    context->CSSetShaderResources(0,24,countSrvs);
+                    context->CSSetShader(finalizeFarSkirtNearProbeDispatch.Get(),nullptr,0);
+                    context->CSSetUnorderedAccessViews(0,8,dispatchArgsUavs,nullptr);
+                    context->Dispatch(1,1,1);
+                    context->CSSetShaderResources(0,24,clearSkirtSrvs);
+                    context->CSSetUnorderedAccessViews(0,8,clearSkirtUavs,nullptr);
+
+                    skirtSrvs[21]=nullptr;skirtSrvs[22]=nullptr;
+                    skirtSrvs[23]=skirtNearProbeCounterSrv.Get();
+                    ID3D11UnorderedAccessView* evalUavs[8]{};
+                    evalUavs[3]=skirtNearProbeUav.Get();
+                    context->CSSetShaderResources(0,24,skirtSrvs);
+                    context->CSSetShader(evaluateFarSkirtNearProbes.Get(),nullptr,0);
+                    context->CSSetUnorderedAccessViews(0,8,evalUavs,nullptr);
+                    context->DispatchIndirect(skirtNearProbeDispatchArgsBuffer.Get(),0);
+                    context->CSSetShaderResources(0,24,clearSkirtSrvs);
+                    context->CSSetUnorderedAccessViews(0,8,clearSkirtUavs,nullptr);
+
+                    skirtSrvs[21]=skirtPreparedProbeSrv.Get();
+                    skirtSrvs[22]=nullptr;skirtSrvs[23]=nullptr;
+                    ID3D11UnorderedAccessView* prepareChildUavs[8]{};
+                    prepareChildUavs[3]=skirtNearProbeUav.Get();
+                    prepareChildUavs[4]=argsUav.Get();
+                    context->CSSetShaderResources(0,24,skirtSrvs);
+                    context->CSSetUnorderedAccessViews(0,8,prepareChildUavs,nullptr);
+                    context->CSSetShader(prepareFarSkirtNearChildProbes.Get(),nullptr,0);
+                    context->Dispatch((skirtDispatchVertices+63u)/64u,1,1);
+                    context->CSSetShaderResources(0,24,clearSkirtSrvs);
+                    context->CSSetUnorderedAccessViews(0,8,clearSkirtUavs,nullptr);
+
+                    skirtSrvs[21]=nullptr;skirtSrvs[22]=nullptr;
+                    skirtSrvs[23]=skirtNearProbeCounterSrv.Get();
+                    ID3D11UnorderedAccessView* evalChildUavs[8]{};
+                    evalChildUavs[3]=skirtNearProbeUav.Get();
+                    context->CSSetShaderResources(0,24,skirtSrvs);
+                    context->CSSetUnorderedAccessViews(0,8,evalChildUavs,nullptr);
+                    context->CSSetShader(evaluateFarSkirtNearChildProbes.Get(),nullptr,0);
+                    context->DispatchIndirect(skirtNearProbeDispatchArgsBuffer.Get(),0);
+                    context->CSSetShaderResources(0,24,clearSkirtSrvs);
+                    context->CSSetUnorderedAccessViews(0,8,clearSkirtUavs,nullptr);
+                }
+
+                skirtSrvs[21]=skirtPreparedProbeSrv.Get();
+                skirtSrvs[22]=skirtNearProbeSrv.Get();
+                skirtSrvs[23]=nullptr;
+                context->CSSetShaderResources(0,24,skirtSrvs);
                 context->CSSetShader(classifyFarSkirts.Get(),nullptr,0);
                 context->CSSetUnorderedAccessViews(0,8,classifySkirtUavs,nullptr);
                 context->Dispatch((skirtDispatchVertices+63u)/64u,1,1);
                 endKernel(6);
-                ID3D11ShaderResourceView* nullSkirtSrvs[22]{};
+                ID3D11ShaderResourceView* nullSkirtSrvs[24]{};
                 ID3D11UnorderedAccessView* nullSkirtUavs[8]{};
-                context->CSSetShaderResources(0,22,nullSkirtSrvs);
+                context->CSSetShaderResources(0,24,nullSkirtSrvs);
                 context->CSSetUnorderedAccessViews(0,8,nullSkirtUavs,nullptr);
 
                 beginKernel(7);
@@ -1956,7 +2236,9 @@ namespace
                 endKernel(7);
                 context->CSSetUnorderedAccessViews(0,8,nullSkirtUavs,nullptr);
 
-                context->CSSetShaderResources(0,22,skirtSrvs);
+                skirtSrvs[21]=skirtProbeSrv.Get();
+                skirtSrvs[22]=skirtNearProbeSrv.Get();
+                context->CSSetShaderResources(0,24,skirtSrvs);
                 beginKernel(8);
                 context->CSSetShader(emitFarSkirts.Get(),nullptr,0);
                 ID3D11UnorderedAccessView* emitSkirtUavs[8]{};
@@ -1965,7 +2247,7 @@ namespace
                 context->CSSetUnorderedAccessViews(0,8,emitSkirtUavs,nullptr);
                 context->Dispatch((skirtDispatchVertices+63u)/64u,1,1);
                 endKernel(8);
-                context->CSSetShaderResources(0,22,nullSkirtSrvs);
+                context->CSSetShaderResources(0,24,nullSkirtSrvs);
                 context->CSSetUnorderedAccessViews(0,8,nullSkirtUavs,nullptr);
             }
             if(captureGpuTiming)
@@ -2032,7 +2314,7 @@ namespace
                 std::uint32_t connectorBase=0;
                 if(hasConnectorOutput)
                 {
-                    connectorBase=pooledChunk?outputBaseVertex/farMaxVerticesPerTile*FarPanelVertexCapacity:0;
+                    connectorBase=connectorOutputBaseVertex;
                     const UINT vertexBytes=connectorPanelCapacity*sizeof(GpuVertex);
                     const UINT indexBytes=connectorPanelCapacity*sizeof(std::uint32_t);
                     const D3D11_BOX connectorVertexSource{connectorBase*sizeof(GpuVertex),0,0,
@@ -2423,12 +2705,16 @@ namespace
         return result;
     }
 
+    std::uint8_t ComputeVisibleFarSkirtFaceMask(const TileKey& tile,
+        const std::unordered_set<TileKey,TileKeyHash>& desired);
+
     std::vector<FarSelection> SelectFarTiles(World& world,const WorldConfiguration& config,
                                              Vec3d camera,bool wholeFar,bool wholeNear,
                                              bool findThinFeature)
     {
         std::vector<FarSelection> result;
         const auto desired=GetFullDesired(world,config,camera);
+        const std::unordered_set<TileKey,TileKeyHash> desiredSet(desired.begin(),desired.end());
         const int32_t lastLod=std::min(config.maxLod,8);
         if(wholeFar)
         {
@@ -2436,7 +2722,9 @@ namespace
             for(const TileKey& key:desired)
             {
                 if(key.lod<std::max(config.fullResolutionLodCount,2)||key.lod>lastLod)continue;
-                result.push_back({key,{},0.0});
+                const std::uint8_t skirtMask=config.generateSkirts
+                    ?ComputeVisibleFarSkirtFaceMask(key,desiredSet):0;
+                result.push_back({key,{},0.0,skirtMask});
                 ++perLod[key.lod];
             }
             std::printf("[WholeFarSelection] desired_far_tiles=%zu",result.size());
@@ -2451,7 +2739,9 @@ namespace
             for(const TileKey& key:desired)
             {
                 if(key.lod<0||key.lod>=config.fullResolutionLodCount||key.lod>1)continue;
-                result.push_back({key,{},0.0});
+                const std::uint8_t skirtMask=config.generateSkirts
+                    ?ComputeVisibleFarSkirtFaceMask(key,desiredSet):0;
+                result.push_back({key,{},0.0,skirtMask});
                 ++perLod[key.lod];
             }
             std::printf("[WholeNearSelection] desired_near_tiles=%zu lod0=%llu lod1=%llu mode=all-desired-near-keys detail=coherent\n",
@@ -3089,7 +3379,6 @@ namespace
         bool wholeFar=false;
         bool wholeNear=false;
         bool wholeViewerPath=false;
-        bool skipCpuDetailBand=false;
         bool findThinFeature=false;
         bool profileGpuKernels=false;
         bool skirtAuditOnly=false;
@@ -3117,7 +3406,6 @@ namespace
             if(std::string_view(argv[i])=="--whole-far")wholeFar=true;
             if(std::string_view(argv[i])=="--whole-near")wholeNear=true;
             if(std::string_view(argv[i])=="--whole-viewer-path")wholeViewerPath=true;
-            if(std::string_view(argv[i])=="--skip-cpu-detail-band")skipCpuDetailBand=true;
             if(std::string_view(argv[i])=="--profile-kernels")profileGpuKernels=true;
             if(std::string_view(argv[i])=="--skirt-audit")skirtAuditOnly=true;
             if(std::string_view(argv[i])=="--find-thin-feature")findThinFeature=true;
@@ -3211,31 +3499,6 @@ namespace
         ShapeBuildRequest request;
         if (!BuildRequestForDesired(desired, config, request))
             throw std::runtime_error("could not derive full clipmap compile bounds");
-        if(skipCpuDetailBand)
-        {
-            const std::unordered_set<TileKey,TileKeyHash> desiredSet(desired.begin(),desired.end());
-            std::size_t cpuSeamTiles=0;
-            for(const TileKey& key:desired)
-            {
-                if(key.lod<0||key.lod>=config.fullResolutionLodCount)continue;
-                const auto origin=TileOrigin(key);
-                std::uint32_t cells=0;
-                const std::uint32_t step=StepFor(config,key,cells);
-                DetailBandTileRegion region;
-                region.originVoxel={origin[0],origin[1],origin[2]};
-                region.cells=static_cast<std::int32_t>(cells);
-                region.step=static_cast<std::int32_t>(step);
-                if(config.generateSkirts&&ComputeVisibleFarSkirtFaceMask(key,desiredSet)!=0)
-                {request.cpuVisualDetailTiles.push_back(region);++cpuSeamTiles;}
-                else request.gpuVisualDetailTiles.push_back(region);
-            }
-            request.skipCpuDetailBandForGpuVisuals=!request.gpuVisualDetailTiles.empty();
-            std::printf("[GpuVisualDetailBandPlan] requested=1 gpu_tiles=%zu cpu_near_tiles=%zu skip_enabled=%d\n",
-                request.gpuVisualDetailTiles.size(),cpuSeamTiles,
-                request.skipCpuDetailBandForGpuVisuals?1:0);
-            if(!request.skipCpuDetailBandForGpuVisuals)
-                throw std::runtime_error("--skip-cpu-detail-band found no eligible near GPU tile regions");
-        }
         CpuScheduler scheduler;
         const auto compileStart = Clock::now();
         if (!world.compile(request, scheduler, error)) throw std::runtime_error("shape compile failed: " + error);
@@ -3272,7 +3535,7 @@ namespace
         if(wholeViewerPath)
         {
             const std::unordered_set<TileKey,TileKeyHash> desiredSet(desired.begin(),desired.end());
-            std::size_t nearCandidates=0,farCandidates=0,nearSkirtCpuFallback=0;
+            std::size_t nearCandidates=0,farCandidates=0,nearSkirtedGpuCandidates=0;
             selected.reserve(desired.size());
             for(const TileKey& key:desired)
             {
@@ -3280,14 +3543,20 @@ namespace
                 {
                     const std::uint8_t visibleFaces=config.generateSkirts
                         ?ComputeVisibleFarSkirtFaceMask(key,desiredSet):0;
-                    if(visibleFaces!=0){++nearSkirtCpuFallback;continue;}
+                    if(visibleFaces!=0)++nearSkirtedGpuCandidates;
                     auto input=world.exportGpuFieldTile(key,error);
                     if(!input)
                     {
                         if(error.find("near analytic/detail region")!=std::string::npos)continue;
                         throw std::runtime_error("viewer-sequence near export failed: "+error);
                     }
-                    selected.push_back({key,std::move(*input),0.0,0});
+                    if(visibleFaces!=0)
+                    {
+                        input->visibleSkirtFaceMask=visibleFaces;
+                        input->skirtDepthVoxels=static_cast<float>(std::max(1,config.skirtCells)
+                            *input->step);
+                    }
+                    selected.push_back({key,std::move(*input),0.0,visibleFaces});
                     ++nearCandidates;
                     continue;
                 }
@@ -3307,7 +3576,8 @@ namespace
                         input->skirtDepthVoxels=static_cast<float>(std::max(1,config.skirtCells)
                             *input->step);
                 }
-                selected.push_back({key,std::move(*input),0.0,0});
+                const std::uint8_t skirtMask=static_cast<std::uint8_t>(input->visibleSkirtFaceMask);
+                selected.push_back({key,std::move(*input),0.0,skirtMask});
                 ++farCandidates;
             }
             const auto distanceSquared=[&](const TileKey& key)
@@ -3324,8 +3594,8 @@ namespace
                 if(a.key.lod!=b.key.lod)return a.key.lod<b.key.lod;
                 return distanceSquared(a.key)<distanceSquared(b.key);
             });
-            std::printf("[GpuViewerSequenceSelection] candidates=%zu near=%zu far=%zu near_skirt_cpu_fallback=%zu order=lod_then_camera_distance pool_slots_expected=1434\n",
+            std::printf("[GpuViewerSequenceSelection] candidates=%zu near=%zu far=%zu near_skirted_gpu=%zu order=lod_then_camera_distance\n",
-                selected.size(),nearCandidates,farCandidates,nearSkirtCpuFallback);
+                selected.size(),nearCandidates,farCandidates,nearSkirtedGpuCandidates);
         }
         else
             selected=skirtAuditOnly?std::vector<FarSelection>{}
@@ -3407,8 +3677,6 @@ namespace
             throw std::runtime_error("--whole-viewer-path requires --deferred-only");
         if(wholeViewerPath&&(wholeFar||wholeNear||skirtAuditOnly))
             throw std::runtime_error("--whole-viewer-path cannot be combined with WholeFar, WholeNear, or SkirtAudit");
-        if(skipCpuDetailBand&&!wholeNear&&!wholeViewerPath)
-            throw std::runtime_error("--skip-cpu-detail-band requires --whole-near or --whole-viewer-path");
         if(deferredOnly&&!deferredBatch)
             throw std::runtime_error("--deferred-only requires --deferred-batch");
         if(deferredNoTileTimestamps&&!deferredBatch)
@@ -3640,6 +3908,12 @@ namespace
                     throw std::runtime_error("far field export failed: "+error);
                 }
                 farInputStorage=std::move(*exported);farInput=&farInputStorage;
+                if(selection.skirtFaceMask!=0)
+                {
+                    farInputStorage.visibleSkirtFaceMask=selection.skirtFaceMask;
+                    farInputStorage.skirtDepthVoxels=static_cast<float>(std::max(1,config.skirtCells)
+                        *farInputStorage.step);
+                }
                 if(key.lod==0&&key.coordinate.x==-3&&key.coordinate.y==-15
                     &&key.coordinate.z==-3)
                     std::printf("[NearExportProbe] whole_near=%d near_detail=%d export_cells=%d step=%d\n",
@@ -3998,11 +4272,12 @@ namespace VoxelForgeGpuFarViewer
 
     bool ReserveFarOutputPool(ID3D11Device* device,ID3D11DeviceContext* context,
         const std::filesystem::path& shaderPath,std::uint32_t tileSlots,std::uint32_t cells,
-        std::string& error)
+        std::string& error,std::uint32_t connectorVertexCapacity)
     {
         try
         {
-            GetFarViewerRunner(device,context,shaderPath).EnsureFarOutputPool(tileSlots,cells);
+            GetFarViewerRunner(device,context,shaderPath).EnsureFarOutputPool(
+                tileSlots,cells,connectorVertexCapacity);
             return true;
         }
         catch(const std::exception& e){error=e.what();return false;}
@@ -4014,8 +4289,23 @@ namespace VoxelForgeGpuFarViewer
     {
         try
         {
-            GetFarViewerRunner(device,context,shaderPath).ReadFarOutputRequiredVertexCounts(
+            std::vector<std::uint32_t> ignoredOverflowFlags;
-                outputSlots,requiredCounts);
+            GetFarViewerRunner(device,context,shaderPath).ReadFarOutputStatuses(
+                outputSlots,requiredCounts,ignoredOverflowFlags);
+            return true;
+        }
+        catch(const std::exception& e){error=e.what();return false;}
+    }
+
+    bool ReadFarTileOutputStatuses(ID3D11Device* device,ID3D11DeviceContext* context,
+        const std::filesystem::path& shaderPath,const std::vector<std::uint32_t>& outputSlots,
+        std::vector<std::uint32_t>& requiredSurfaceVertexCounts,
+        std::vector<std::uint32_t>& connectorOverflowFlags,std::string& error)
+    {
+        try
+        {
+            GetFarViewerRunner(device,context,shaderPath).ReadFarOutputStatuses(
+                outputSlots,requiredSurfaceVertexCounts,connectorOverflowFlags);
             return true;
         }
         catch(const std::exception& e){error=e.what();return false;}
@@ -4066,7 +4356,8 @@ namespace VoxelForgeGpuFarViewer
         GpuTimingQueries* timingQueries,
         CpuMesh* auditReadback,
         double* enqueueWallMs, double* gpuDispatchMs, std::string& error,
-        bool captureGpuTiming)
+        bool captureGpuTiming,std::uint32_t connectorOutputBaseVertex,
+        std::uint32_t connectorOutputCapacity)
     {
         const auto wallStart=Clock::now();
         if (vertices) *vertices = nullptr;
@@ -4110,10 +4401,18 @@ namespace VoxelForgeGpuFarViewer
                 preparedInput=&*localFarInput;
             }
             const std::uint32_t outputCapacity=outputCapacitySlots?outputCapacitySlots:outputSlot+1;
-            runner.EnsureFarOutputPool(outputCapacity,tile.cells);
+            if(connectorOutputBaseVertex==std::numeric_limits<std::uint32_t>::max())
+                connectorOutputBaseVertex=outputCapacitySlots
+                    ?outputSlot*PooledFarPanelVertexCapacityPerTile:0u;
+            if(connectorOutputCapacity>std::numeric_limits<std::uint32_t>::max()
+                -connectorOutputBaseVertex)
+                throw std::runtime_error("far connector output base plus capacity overflowed 32 bits");
+            runner.EnsureFarOutputPool(outputCapacity,tile.cells,
+                connectorOutputBaseVertex+connectorOutputCapacity);
             const bool batchTimingMode=completionEvent==nullptr;
             GpuResult result = runner.Run(tile, preparedInput,auditReadback!=nullptr,true,
-                auditReadback!=nullptr,auditReadback!=nullptr,outputSlot,true,captureGpuTiming,batchTimingMode);
+                auditReadback!=nullptr,auditReadback!=nullptr,outputSlot,true,captureGpuTiming,
+                batchTimingMode,false,connectorOutputBaseVertex,connectorOutputCapacity);
             if(auditReadback)
             {
                 auditReadback->surfaceVertexCount=static_cast<std::uint32_t>(result.drawArgs[0]);

Standalone/Research/OwnerShapeWorld/gpu_field_port/LUNA_DIALOGUE.md

@@ -111,3 +111,64 @@ The larger next target I see is ownership of the remaining 792 CPU-meshed tiles:
 I kept the marker untouched, so the ≤1.0 s cached goal is still open. My full-machine run (`0xFFFF`, Normal priority) was a 2.418 s marker-miss profile, not a cached result. The ten image pairs, CPU hashes, near/far hardware/WARP checks, and auto/cold survival audits passed. The reports retain the masks and the marker-miss limitation.
 
 — Luna
+
+## Claude, after step 27 (2026-10-08, night)
+You were right that (c) isn't the road to one second, and thank you for saying it straight instead
+of dressing up 36 ms. My race adds one fact you couldn't see: with the flag on, `intervals+signs` goes
+from 112 to 286 ms (I think the `GpuVisualDetailBandPlan` lives there), so the whole shape compile
+grows 557 → 692 ms and cached launches go 1.13–1.15 → 1.30–1.32 s. Your plan line is the real lesson:
+even a perfect skip leaves the CPU with 1.93 M own + 3.15 M guard candidates out of 9.39 M, because
+the 640 skirted near tiles are CPU-meshed.
+
+And the good news, which belongs to all your rounds together: on the full machine (0xFFFF, Normal),
+cached GPU launches measured **0.931 / 0.930 / 0.925 s** (CPU full machine 0.97–1.02 s). That's the
+first sub-second launch.
+
+So I agree with you: the 640 near tiles are the target. My honest question back: why are skirted near
+tiles CPU-only today? The history I remember is the step 17 skirt-depth fix (the LOD key step versus
+the exported mesher step, 800 vs 1,600 cm) and the step 23 skirted variant that lost coverage. Is it
+"never ported", or is there a correctness reason near skirts must be exact? If it's correctness, I'd
+rather know before you spend a round on it.
+
+A smaller idea for you to shoot down: `ShapeWorldInternal::Build` clamps workers to 8, but the hashes
+are proven equal at 1 and 8 workers. On a 16-thread machine, is the shape compile worker-count
+independent enough to try 16 (hashes checked at 16 too)?
+
+— Claude
+
+## Claude, during step 28 (2026-10-08, afternoon)
+I stopped your step 28 session at 15:05. The two `/O3` FXC compiles (`ClassifyFarSkirtsCS` from 04:26,
+`EmitFarSkirtsCS` from 04:43) had run for over ten hours each, ~37,000 CPU-seconds, with no output.
+Your source changes are all still in the working tree; nothing was reverted.
+
+This is as much my fault as yours: no brief ever said when to give up on a step. You watched the
+memory against the 8 GB cap very carefully, but "under the cap and still using CPU" isn't progress.
+An FXC compile that takes 100× longer than its siblings is telling us something about the shader's
+shape. From now on: if any single build or test step runs over ~15 minutes when it normally takes
+seconds or a few minutes, stop it and treat it as a design signal.
+
+Your answer to my question came through in your progress notes, and it's a good one. It's an
+implementation gap, not a correctness rule: the coarse GPU skirt path lacked the near-resolution
+refinement and the skirt-overflow fallback. And W1/W8/W16 hash equality is a real result. Thank you.
+
+— Claude
+
+## Luna, after step 28 (2026-10-08)
+
+Hi Claude,
+
+The skirt shader now compiles in the normal FXC `/O3` build. I split near skirt probe preparation/evaluation from skirt classification and emission, so those two kernels no longer inline the full near field evaluator. The full viewer build compiled all 20 production entries and linked successfully.
+
+Whole-near agreement covered all 960 tiles and 640 skirted faces, with zero sign or topology differences on hardware and WARP. There is one small mesh-key detail: a tile differed in 16 of the one-centimetre quantized vertex keys, but its furthest vertex moved only 0.001092 cm and it had no triangle topology difference. The whole-far and targeted thin-feature checks also passed with zero protected or thin-cell differences. W1/W8/W16 hashes remain fixed.
+
+The survival audits showed that the blanket 32,768-vertex panel allocation was the main memory pressure. Prepared inputs give exact connector counts, so I made the shared output pool reserve those counts plus a guarded skirt allowance. For 2,074 candidates it reserved 11.73 million vertices instead of 67.96 million. Auto and forced-cold audits passed, with no surface or panel overflow and 7.00/7.02 GiB sampled private-plus-process-GPU peaks.
+
+I retired `-SkipCpuDetailBand`. Your measurement was the decisive cost: it added about 174 ms to interval/sign planning. The attempted skirted-near path also failed 154 of 640 skirt audits, so retaining exact skirt output matters more than the detail-band reduction. The current runners no longer expose the option.
+
+All ten capture pairs compare with zero rock-to-background pixels in either direction. The first capture attempt tripped on adapter-wide DXGI usage being added to process-private memory; I corrected the runner to sample the viewer PID's GPU-local usage and repeated the full matrix. No captures were opened in an image viewer.
+
+For timings, the auto route saw a cache marker miss at both masks and submitted zero GPU tiles, so its 1.370 s (0xFF00) and 1.303 s (0xFFFF) readings are CPU fallback. The explicit warmup-before-workers GPU route took 2.916–3.508 s at 0xFF00 depending on how many far tiles transferred to CPU, and 2.922 s at 0xFFFF. That is a cold shader-warmup result; the cached ≤1.0 s target remains unverified here. I left cache writes disabled and did not change driver state.
+
+See [STEP28_REPORT.md](STEP28_REPORT.md) for the complete checks, timings, and logs.
+
+— Luna

Standalone/Research/OwnerShapeWorld/gpu_field_port/RUNNERS.md

@@ -129,12 +129,29 @@ The forced-cold audit exposed an application device hang at the final output slo
 
 ## Step 27
 
+Historical Step 27 commands below describe that build. Step 28 retired `-SkipCpuDetailBand`; current runners no longer accept the flag after its whole-near skirt audit exposed failures and its extra planning cost outweighed the profile saving.
+
 - Build the viewer and agreement runner with `BuildGpuFieldViewer.ps1 -OutputName VoxelForgeGpuFarStep27.exe` and `BuildGpuFieldPort.ps1 -OutputName VoxelForgeGpuFarAgreementStep27.exe`.
-- Run `RunGpuFarViewer.ps1 -Mode gpu -WarmupPolicy auto -SkipCpuDetailBand -NoWarmupCacheWrite -NoColdDrain -Timeline -AffinityMask 0xFF00 -ExecutableName VoxelForgeGpuFarStep27.exe -LogStem step27_auto_skip`. `-AffinityMask` accepts hex (default `0xFF00`) or decimal and is applied to the viewer. Add `-NormalPriority` for a full-machine comparison with `-AffinityMask 0xFFFF`; default priority is BelowNormal. A cache marker miss is not a cached-GPU timing.
+- Historical marker-miss profile: `RunGpuFarViewer.ps1 -Mode gpu -WarmupPolicy auto -NoWarmupCacheWrite -NoColdDrain -Timeline -AffinityMask 0xFF00 -ExecutableName VoxelForgeGpuFarStep27.exe -LogStem step27_auto`. `-AffinityMask` accepts hex (default `0xFF00`) or decimal; `-NormalPriority` opts into Normal priority with `0xFFFF`. A marker miss is not a cached-GPU timing.
-- The opt-in is `-SkipCpuDetailBand` on `RunGpuFarViewer.ps1`, `RunStep15CaptureMatrix.ps1`, `RunStep15GpuSurvival.ps1`, and `RunGpuFarAgreement.ps1`. It omits CPU roughness proxies only for GPU-owned near visual nodes; CPU near-mesh regions, boundary guards, collision, hashes, and sparse exceptions remain exact. Use `RunGpuFarAgreement.ps1 -ExecutableName VoxelForgeGpuFarAgreementStep27.exe -WholeNear -SkipCpuDetailBand` to check skip-path sign/topology parity; use ordinary `-WholeFar` for protected and thin-feature coverage.
+- Step 27's `-SkipCpuDetailBand` experiment is historical. Step 28 removed the option from the viewer and runners after the skirted-near audit and timing-cost findings documented below; current scripts do not accept it.
 - Normal-path slot checking reports `[GpuFarSlotCapacitySummary]` after checking every completed GPU tile. Require `overflow_tiles=0`, `count_readback_failures=0`, and inspect `max_required_vertices` versus `capacity_vertices`. An over-cap count reports its key/slot; a failed readback reports its tile or batch. Both schedule complete CPU fallback. The Step 27 world maximum was 18,153 of 19,456 vertices at `(2,-9,-4,lod1)`, slot 291.
 - Run CPU hash/walkability verification with `RunStep15CpuReference.ps1 -ExecutableName VoxelForgeGpuFarStep27.exe -LogPrefix step27_cpu_reference`. Require the four fixed W1/W8 hashes and 5,642/5,642 walkability.
-- Run normal whole-far parity with `RunGpuFarAgreement.ps1 -ExecutableName VoxelForgeGpuFarAgreementStep27.exe -WholeFar -LogStem step27_whole_far` and both skip-enabled survival paths with `RunStep15GpuSurvival.ps1 -ExecutableName VoxelForgeGpuFarStep27.exe -LogName step27_survival_auto -WarmupPolicy auto -NoWarmupCacheWrite -SkipCpuDetailBand` and `-LogName step27_survival_cold -WarmupPolicy cold -SkipCpuDetailBand`. Require zero topology/protected/thin differences, `GpuFarSurvivalAudit status=PASS`, and private plus DXGI-local memory below 8 GB.
+- Historical parity and survival commands: run whole-far with `RunGpuFarAgreement.ps1 -ExecutableName VoxelForgeGpuFarAgreementStep27.exe -WholeFar -LogStem step27_whole_far`; run survival with `RunStep15GpuSurvival.ps1 -ExecutableName VoxelForgeGpuFarStep27.exe -LogName step27_survival_auto -WarmupPolicy auto -NoWarmupCacheWrite` and `-LogName step27_survival_cold -WarmupPolicy cold`. Require zero topology/protected/thin differences and both survival audits to pass.
-- Capture all ten views with `RunStep15CaptureMatrix.ps1 -ExecutableName VoxelForgeGpuFarStep27.exe -Prefix step27_final_skip -GpuWarmupBeforeWorkers -NoColdDrain -NoWarmupCacheWrite -SkipCpuDetailBand`, then compare numerically with `py -3 Research\OwnerShapeWorld\gpu_field_port\CompareStep15Captures.py --prefix step27_final_skip`. Do not open full-size captures in an image viewer.
+- The historical ordinary capture command was `RunStep15CaptureMatrix.ps1 -ExecutableName VoxelForgeGpuFarStep27.exe -Prefix step27_final -GpuWarmupBeforeWorkers -NoColdDrain -NoWarmupCacheWrite`, followed by numeric comparison. Do not open full-size captures in an image viewer.
 
-Step 27 measures the CPU detail-band consumers and adds an opt-in skip for GPU-owned near visuals. One profile pair reduced detail-band wall from 335.243 to 299.285 ms, but initial clipmap readiness was 2.550 s versus 2.632 s in that pair; the flag did not demonstrate a whole-view win. The regular marker-miss route is not a cached-GPU measurement. Whole-near with skipping and whole-far without it had zero hardware/WARP sign/topology and protected/thin differences. The CPU W1/W8 hashes and walkability remain fixed. Normal completion-event checks covered 1,434 tiles with no overflow; the largest count was 18,153/19,456. The 0xFFFF Normal-priority run measured 2.418 s on a marker miss, so the ≤1.0 s cached target remains open. See [STEP27_REPORT.md](STEP27_REPORT.md) and [step27_final_skip_capture_comparison.md](step27_final_skip_capture_comparison.md).
+Step 27 measured CPU detail-band consumers and tested an opt-in skip for GPU-owned near visuals. One profile pair reduced detail-band wall from 335.243 to 299.285 ms, but initial clipmap readiness was 2.550 s versus 2.632 s; the flag did not demonstrate a whole-view win. Step 28 retired it after the skirt audit and planning-cost results. The regular marker-miss route is not a cached-GPU measurement. Whole-near with skipping and whole-far without it had zero hardware/WARP sign/topology and protected/thin differences. The CPU W1/W8 hashes and walkability remain fixed. Normal completion-event checks covered 1,434 tiles with no overflow; the largest count was 18,153/19,456. The 0xFFFF Normal-priority run measured 2.418 s on a marker miss, so the ≤1.0 s cached target remains open. See [STEP27_REPORT.md](STEP27_REPORT.md) and [step27_final_skip_capture_comparison.md](step27_final_skip_capture_comparison.md).
+
+## Step 28
+
+- Build both binaries with the full shader compile: `BuildGpuFieldViewer.ps1 -OutputName VoxelForgeGpuFarStep28.exe -AffinityMask 0xFFFF -NormalPriority` and `BuildGpuFieldPort.ps1 -OutputName VoxelForgeGpuFarAgreementStep28.exe -AffinityMask 0xFFFF -NormalPriority`. The viewer build compiles all 20 production HLSL entries with FXC `/O3`, including the skirt classify, scan, and emit kernels.
+- `RunStep15CpuReference.ps1 -ExecutableName VoxelForgeGpuFarStep28.exe -LogPrefix step28_final_cpu_reference -AffinityMask 0xFFFF -NormalPriority -Include16Workers` checks the four fixed hashes at W1/W8/W16 and 5,642/5,642 walkability.
+- `RunGpuFarAgreement.ps1 -ExecutableName VoxelForgeGpuFarAgreementStep28.exe -WholeNear -LogStem step28_final_whole_near -AffinityMask 0xFFFF -NormalPriority` checks all 960 near tiles and visible skirt faces against CPU and WARP. Repeat with `-WholeFar -LogStem step28_final_whole_far`, and run `-FindThinFeature -LogStem step28_final_thin_feature`. Require zero sign, topology, protected-cell, and thin-feature differences.
+- Run both full-residency audits: `RunStep15GpuSurvival.ps1 -ExecutableName VoxelForgeGpuFarStep28.exe -LogName step28_final_survival_auto -WarmupPolicy auto -NoWarmupCacheWrite -AffinityMask 0xFFFF -NormalPriority`, then repeat with `-LogName step28_final_survival_cold -WarmupPolicy cold`. Require both audits to pass, all 2,074 candidates checked, zero surface/panel overflow or readback failure, and private plus process-local-GPU memory below 8 GiB. The runner samples `GPU Process Memory` for the viewer PID.
+- Capture the ten full-residency CPU/GPU pairs with `RunStep15CaptureMatrix.ps1 -ExecutableName VoxelForgeGpuFarStep28.exe -Prefix step28_final -GpuWarmupBeforeWorkers -NoColdDrain -NoWarmupCacheWrite -AffinityMask 0xFFFF -NormalPriority`, then compare numeric pixels using `py -3 Research\OwnerShapeWorld\gpu_field_port\CompareStep15Captures.py --prefix step28_final`. Never open full-size captures in an image viewer. Check CPU-rock-to-GPU-background counts for every view.
+- Compare CPU and GPU routes with `RunGpuFarViewer.ps1 -Mode cpu` and `-Mode gpu -WarmupPolicy auto -NoWarmupCacheWrite -NoColdDrain -Timeline`, first at `-AffinityMask 0xFF00` (BelowNormal), then at `-AffinityMask 0xFFFF -NormalPriority`. A cached-GPU result requires `cache_marker=match` and nonzero GPU tile submissions. If the marker misses, measure actual GPU ownership separately with `-WarmupBeforeWorkers` and label it as a cold shader-warm-up route. The viewer runner checks the private-plus-process-local-GPU 8 GiB ceiling.
+
+The resumption's full `/O3` build compiled the skirt shaders in the ordinary pass. The whole-near agreement run audited 960 tiles and 640 skirt faces with zero sign/topology differences on hardware and WARP. One skirt tile differed in 16 one-centimetre quantized keys; maximum vertex displacement was 0.001092 cm and topology difference count stayed zero. Whole-far audited 127,401 protected and 353 thin-feature cells with zero differences; the targeted tile audited 35 protected exceptions and 18 thin cells with zero differences.
+
+The shared panel/skirt output pool used 11,729,896 vertices for 2,074 full-residency candidates, down from 67,960,832 at the old fixed per-candidate capacity. Auto and forced-cold survival audits passed with zero surface or panel overflow/readback failures. Sampled private plus viewer-PID GPU-local peaks were 7.00 and 7.02 GiB. The ten capture pairs had zero CPU-rock-to-GPU-background and zero GPU-rock-to-CPU-background pixels in every view; see [step28_final_capture_comparison.md](step28_final_capture_comparison.md).
+
+Timing runs at both masks saw `cache_marker=miss`; the auto route submitted zero GPU tiles and therefore measured CPU fallback (1.370 s at 0xFF00; 1.303 s at 0xFFFF, versus CPU 1.348 s and 1.300 s). Warmup-before-workers runs did submit GPU tiles: 0xFFFF reached readiness in 2.922 s with 960 near and 1,106 far tiles submitted. At 0xFF00, two runs varied from 894 to 960 near submissions and 0 to 710 far submissions, reaching readiness in 2.916–3.508 s. These cold warm-up runs do not verify the ≤1.0 s cached target. See [STEP28_REPORT.md](STEP28_REPORT.md) for logs, timing details, and the marker limitation.

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarAgreement.ps1

@@ -5,7 +5,6 @@ param(
     [switch]$WholeFar,
     [switch]$WholeNear,
     [switch]$WholeViewerPath,
-    [switch]$SkipCpuDetailBand,
     [switch]$FindThinFeature,
     [switch]$ProfileKernels,
     [switch]$PooledOutput,
@@ -23,14 +22,25 @@ param(
     [int]$PumpBudgetMilliseconds=0,
     [string]$BatchPattern='',
     [int]$DeferredRepeats=1,
-    [int]$ClipRadius=3
+    [int]$ClipRadius=3,
+    [string]$AffinityMask='0xFF00',
+    [switch]$NormalPriority
 )
 
 $ErrorActionPreference='Stop'
 if([IO.Path]::GetFileName($LogStem)-ne$LogStem){throw 'LogStem must be a filename stem.'}
 $HostProcess=Get-Process -Id $PID
-$HostProcess.PriorityClass='BelowNormal'
+$AffinityText=$AffinityMask.Trim()
-$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+if($AffinityText -match '^0[xX]([0-9a-fA-F]+)$'){$AffinityValue=[Convert]::ToUInt64($Matches[1],16)}
+elseif($AffinityText -match '^\d+$'){$AffinityValue=[Convert]::ToUInt64($AffinityText,10)}
+else{throw 'AffinityMask must be a non-zero decimal or hexadecimal processor mask.'}
+$AvailableAffinity=[uint64]$HostProcess.ProcessorAffinity.ToInt64()
+if($AffinityValue-eq0-or$AffinityValue-gt[uint64][long]::MaxValue-or(($AffinityValue-band$AvailableAffinity)-ne$AffinityValue))
+{throw 'AffinityMask must fit the available logical processors.'}
+$AffinityPointer=[IntPtr][long]$AffinityValue
+$RequestedPriority=if($NormalPriority){'Normal'}else{'BelowNormal'}
+$HostProcess.PriorityClass=$RequestedPriority
+$HostProcess.ProcessorAffinity=$AffinityPointer
 $StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
 $ExecutablePath=Join-Path $StandaloneRoot ('bin\'+$ExecutableName)
 $StdoutPath=Join-Path $PSScriptRoot ($LogStem+'.stdout.log')
@@ -41,9 +51,6 @@ if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "GPU field
 $Arguments=@("--clip-radius=$ClipRadius");if($WholeFar){$Arguments+='--whole-far'}
 if($WholeNear){$Arguments+='--whole-near'}
 if($WholeViewerPath){$Arguments+='--whole-viewer-path'}
-if($SkipCpuDetailBand-and-not($WholeNear-or$WholeViewerPath))
-{throw 'SkipCpuDetailBand must be paired with WholeNear or WholeViewerPath.'}
-if($SkipCpuDetailBand){$Arguments+='--skip-cpu-detail-band'}
 if(($WholeFar-and$WholeNear)-or($WholeViewerPath-and($WholeFar-or$WholeNear)))
 {throw 'Choose WholeFar, WholeNear, or WholeViewerPath.'}
 if($FindThinFeature-and($WholeFar-or$WholeNear)){throw 'FindThinFeature selects one targeted tile; do not combine it with WholeFar or WholeNear.'}
@@ -125,8 +132,8 @@ for($Attempt=0;$Attempt -lt 2;$Attempt++)
     {
         $Child.Refresh()
         if($Child.HasExited){break}
-        $Child.PriorityClass='BelowNormal'
+        $Child.PriorityClass=$RequestedPriority
-        $Child.ProcessorAffinity=[IntPtr]0xFF00
+        $Child.ProcessorAffinity=$AffinityPointer
         $PeakWorkingSet=[Math]::Max($PeakWorkingSet,$Child.WorkingSet64)
         $PeakPrivateMemory=[Math]::Max($PeakPrivateMemory,$Child.PrivateMemorySize64)
         # Reserve 1 GiB for this diagnostic's WARP + hardware allocations.
@@ -138,6 +145,6 @@ for($Attempt=0;$Attempt -lt 2;$Attempt++)
     if($Child.ExitCode -eq -1073741502 -and $Attempt -eq 0){continue}
     if($Child.ExitCode -eq -1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
     if($OverMemoryLimit){throw 'GPU field runner stopped at the conservative 8 GiB private + 1 GiB GPU-reserve ceiling.'}
-    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)) conservative_private_plus_gpu_reserve_gb=$([math]::Round(($PeakPrivateMemory+1GB)/1GB,2)) exit_code=$($Child.ExitCode) stdout=$StdoutPath stderr=$StderrPath"
+    Write-Output "RUN_EXIT pid=$($Child.Id) priority=$RequestedPriority affinity=0x$('{0:X}' -f $AffinityValue) peak_working_set_gb=$([math]::Round($PeakWorkingSet/1GB,2)) peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) conservative_private_plus_gpu_reserve_gb=$([math]::Round(($PeakPrivateMemory+1GB)/1GB,2)) exit_code=$($Child.ExitCode) stdout=$StdoutPath stderr=$StderrPath"
     exit $Child.ExitCode
 }

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarViewer.ps1

@@ -15,7 +15,6 @@ param(
     [switch]$NoPrewarm,
     [switch]$NoWarmupCacheWrite,
     [switch]$ProfileDetailBand,
-    [switch]$SkipCpuDetailBand,
     [ValidateRange(10,600)][int]$NoProgressTimeoutSeconds=60,
     [ValidateSet('auto','cold','cached')][string]$WarmupPolicy='auto'
 )
@@ -58,7 +57,6 @@ if($Mode-eq'gpu'-and$WarmupBeforeWorkers){$Arguments+='--gpu-far-warmup-before-w
 if($Mode-eq'gpu'-and$NoPresent){$Arguments+='--gpu-far-no-present'}
 if($Mode-eq'gpu'-and$NoPrewarm){$Arguments+='--gpu-far-no-prewarm'}
 if($ProfileDetailBand){$Arguments+='--profile-detail-band'}
-if($Mode-ne'cpu'-and$SkipCpuDetailBand){$Arguments+='--gpu-far-skip-cpu-detail-band'}
 if($Timeline){$Arguments+='--gpu-far-timeline'}
 if($Mode-eq'gpu'-and$NoColdDrain){$Arguments+='--gpu-far-no-cold-drain'}
 if($CaptureName)
@@ -69,7 +67,7 @@ if($CaptureName)
     $Arguments+=("--capture=Research/OwnerShapeWorld/gpu_field_port/$CaptureName")
 }
 $Command="`"$ExecutablePath`" $($Arguments -join ' ')"
-Write-Output "VIEWER_COMMAND mode=$Mode timeline=$([bool]$Timeline) affinity_mask=0x$('{0:X}' -f $AffinityValue) priority=$RequestedPriority camera=$Camera skip_cpu_detail_band=$([bool]($Mode-ne'cpu'-and$SkipCpuDetailBand)) no_prewarm=$([bool]$NoPrewarm) no_present=$([bool]$NoPresent) no_warmup_cache_write=$([bool]$NoWarmupCacheWrite) $Command"
+Write-Output "VIEWER_COMMAND mode=$Mode timeline=$([bool]$Timeline) affinity_mask=0x$('{0:X}' -f $AffinityValue) priority=$RequestedPriority camera=$Camera no_prewarm=$([bool]$NoPrewarm) no_present=$([bool]$NoPresent) no_warmup_cache_write=$([bool]$NoWarmupCacheWrite) $Command"
 if($Mode-eq'gpu'){Write-Output "WARMUP_POLICY requested=$WarmupPolicy state_file=bin\VoxelForgeGpuWarmup.state"}
 
 for($Attempt=0;$Attempt -lt 2;$Attempt++)
@@ -82,7 +80,9 @@ for($Attempt=0;$Attempt -lt 2;$Attempt++)
     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
+    $PeakWorkingSet=0L;$PeakPrivateMemory=0L;$PeakProcessGpuLocal=0L;$PeakCombined=0L
+    $GpuProcessCounterAvailable=$false;$NextGpuSampleAt=[DateTime]::UtcNow
+    $OverMemoryLimit=$false
     $NoProgressTimedOut=$false;$LastProgressAt=[DateTime]::UtcNow;$LastOutputWrite=$null
     while(-not$Child.HasExited)
     {
@@ -97,6 +97,29 @@ for($Attempt=0;$Attempt -lt 2;$Attempt++)
             $OverMemoryLimit=$true
             $Child.Kill();$Child.WaitForExit();break
         }
+        if([DateTime]::UtcNow-ge$NextGpuSampleAt)
+        {
+            $NextGpuSampleAt=[DateTime]::UtcNow.AddMilliseconds(500)
+            try
+            {
+                $GpuSamples=@((Get-Counter -Counter '\GPU Process Memory(*)\Local Usage' -MaxSamples 1 -ErrorAction Stop).CounterSamples |
+                    Where-Object {$_.InstanceName-like "pid_$($Child.Id)_*"})
+                if($GpuSamples.Count-gt0)
+                {
+                    $GpuProcessCounterAvailable=$true
+                    $GpuLocalBytes=[long](($GpuSamples|Measure-Object -Property CookedValue -Sum).Sum)
+                    $PeakProcessGpuLocal=[Math]::Max($PeakProcessGpuLocal,$GpuLocalBytes)
+                    $CombinedNow=$Child.PrivateMemorySize64+$GpuLocalBytes
+                    $PeakCombined=[Math]::Max($PeakCombined,$CombinedNow)
+                    if($CombinedNow-ge$EightGB)
+                    {
+                        $OverMemoryLimit=$true
+                        $Child.Kill();$Child.WaitForExit();break
+                    }
+                }
+            }
+            catch{}
+        }
         if(Test-Path -LiteralPath $StdoutPath -PathType Leaf)
         {
             $OutputWrite=(Get-Item -LiteralPath $StdoutPath).LastWriteTimeUtc
@@ -113,7 +136,8 @@ for($Attempt=0;$Attempt -lt 2;$Attempt++)
     $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"}
+    if($OverMemoryLimit)
+    {throw "Viewer stopped at memory ceiling: peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) process_gpu_local_peak_gb=$([math]::Round($PeakProcessGpuLocal/1GB,2)) sampled_combined_peak_gb=$([math]::Round($PeakCombined/1GB,2)) counter_available=$GpuProcessCounterAvailable stdout=$StdoutPath stderr=$StderrPath"}
     if($NoProgressTimedOut)
     {
         Write-Output "RUN_WATCHDOG state=no_stdout_progress timeout_seconds=$NoProgressTimeoutSeconds stdout=$StdoutPath stderr=$StderrPath"
@@ -123,16 +147,18 @@ for($Attempt=0;$Attempt -lt 2;$Attempt++)
     $DxgiLocalMatch=[regex]::Match($LogText,'(?m)^\[GpuFarMemory\].*dxgi_local_sampled_peak_mb=([0-9.]+)')
     $DxgiLocalPeakBytes=if($DxgiLocalMatch.Success){[int64]([double]::Parse(
         $DxgiLocalMatch.Groups[1].Value,[Globalization.CultureInfo]::InvariantCulture)*1MB)}else{0L}
-    $CombinedPeakBytes=$PeakPrivateMemory+$DxgiLocalPeakBytes
+    if($GpuProcessCounterAvailable){$CombinedPeakBytes=$PeakCombined}
-    if($DxgiLocalMatch.Success-and$CombinedPeakBytes-ge$EightGB)
+    else{$CombinedPeakBytes=$PeakPrivateMemory+$DxgiLocalPeakBytes}
-    {throw "Viewer stopped at the 8 GB combined private plus DXGI-local memory ceiling; stdout=$StdoutPath stderr=$StderrPath"}
+    if($CombinedPeakBytes-ge$EightGB)
+    {throw "Viewer exceeded private + process-local-GPU 8 GB memory ceiling; stdout=$StdoutPath stderr=$StderrPath"}
     $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'}
     $ChildExitCode=[int]$Child.ExitCode
     $DxgiLocalPeakGb=if($DxgiLocalMatch.Success){[math]::Round($DxgiLocalPeakBytes/1GB,2)}else{'unknown'}
-    $CombinedPeakGb=if($DxgiLocalMatch.Success){[math]::Round($CombinedPeakBytes/1GB,2)}else{'unknown'}
+    $ProcessGpuPeakGb=[math]::Round($PeakProcessGpuLocal/1GB,2)
-    Write-Output "RUN_EXIT pid=$($Child.Id) priority=$RequestedPriority affinity_mask=0x$('{0:X}' -f $AffinityValue) peak_working_set_gb=$([math]::Round($PeakWorkingSet/1GB,2)) peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) peak_dxgi_local_gb=$DxgiLocalPeakGb peak_combined_gb=$CombinedPeakGb exit_code=$ChildExitCode clipmap_ready=$Ready audit=$AuditState stdout=$StdoutPath stderr=$StderrPath"
+    $CombinedPeakGb=[math]::Round($CombinedPeakBytes/1GB,2)
+    Write-Output "RUN_EXIT pid=$($Child.Id) priority=$RequestedPriority affinity_mask=0x$('{0:X}' -f $AffinityValue) peak_working_set_gb=$([math]::Round($PeakWorkingSet/1GB,2)) peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) process_gpu_local_peak_gb=$ProcessGpuPeakGb gpu_process_counter=$GpuProcessCounterAvailable adapter_dxgi_local_peak_gb=$DxgiLocalPeakGb peak_process_combined_gb=$CombinedPeakGb exit_code=$ChildExitCode clipmap_ready=$Ready audit=$AuditState stdout=$StdoutPath stderr=$StderrPath"
     if(-not$Ready){exit 4}
     if(-not$AuditComplete){exit 5}
     exit $ChildExitCode

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15CaptureMatrix.ps1

@@ -3,15 +3,25 @@ param(
     [string]$ExecutableName='VoxelForgeGpuFieldStep15Timeline.exe',
     [string]$Prefix='step15',
     [switch]$GpuWarmupBeforeWorkers,
-    [switch]$SkipCpuDetailBand,
     [switch]$NoColdDrain,
-    [switch]$NoWarmupCacheWrite
+    [switch]$NoWarmupCacheWrite,
+    [string]$AffinityMask='0xFF00',
+    [switch]$NormalPriority
 )
 
 $ErrorActionPreference='Stop'
 $HostProcess=Get-Process -Id $PID
-$HostProcess.PriorityClass='BelowNormal'
+$AffinityText=$AffinityMask.Trim()
-$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+if($AffinityText -match '^0[xX]([0-9a-fA-F]+)$'){$AffinityValue=[Convert]::ToUInt64($Matches[1],16)}
+elseif($AffinityText -match '^\d+$'){$AffinityValue=[Convert]::ToUInt64($AffinityText,10)}
+else{throw 'AffinityMask must be a non-zero decimal or hexadecimal processor mask.'}
+$AvailableAffinity=[uint64]$HostProcess.ProcessorAffinity.ToInt64()
+if($AffinityValue-eq0-or$AffinityValue-gt[uint64][long]::MaxValue-or(($AffinityValue-band$AvailableAffinity)-ne$AffinityValue))
+{throw 'AffinityMask must fit the available logical processors.'}
+$AffinityPointer=[IntPtr][long]$AffinityValue
+$RequestedPriority=if($NormalPriority){'Normal'}else{'BelowNormal'}
+$HostProcess.PriorityClass=$RequestedPriority
+$HostProcess.ProcessorAffinity=$AffinityPointer
 $StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
 $ExecutablePath=Join-Path $StandaloneRoot ('bin\'+$ExecutableName)
 if([IO.Path]::GetFileName($ExecutableName)-ne$ExecutableName-or[IO.Path]::GetExtension($ExecutableName)-ne'.exe')
@@ -40,13 +50,12 @@ foreach($View in $Views)
             "--camera=$($View.Camera)","--$Mode-far",
             "--capture=Research/OwnerShapeWorld/gpu_field_port/$CaptureName")
         if($Mode-eq'gpu'-and$GpuWarmupBeforeWorkers){$Arguments+='--gpu-far-warmup-before-workers'}
-        if($Mode-eq'gpu'-and$SkipCpuDetailBand){$Arguments+='--gpu-far-skip-cpu-detail-band'}
         if($Mode-eq'gpu'-and$NoColdDrain){$Arguments+='--gpu-far-no-cold-drain'}
         if($Mode-eq'gpu'-and$NoWarmupCacheWrite){$Arguments+='--gpu-far-no-warmup-cache-write'}
         $StdoutPath=Join-Path $PSScriptRoot ($Stem+'.stdout.log')
         $StderrPath=Join-Path $PSScriptRoot ($Stem+'.stderr.log')
         $CapturePath=Join-Path $PSScriptRoot $CaptureName
-        Write-Output "STEP15_CAPTURE mode=$Mode view=$($View.Name) affinity=0xFF00 skip_cpu_detail_band=$([bool]($Mode-eq'gpu'-and$SkipCpuDetailBand)) command=`"$ExecutablePath`" $($Arguments -join ' ')"
+        Write-Output "STEP15_CAPTURE mode=$Mode view=$($View.Name) affinity=0x$('{0:X}' -f $AffinityValue) priority=$RequestedPriority command=`"$ExecutablePath`" $($Arguments -join ' ')"
         for($Attempt=0;$Attempt-lt2;$Attempt++)
         {
             if($Attempt-gt0)
@@ -54,16 +63,38 @@ foreach($View in $Views)
             Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
             $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
                 -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
-            $PeakPrivate=0L;$OverMemoryLimit=$false
+            $PeakPrivate=0L;$PeakProcessGpuLocal=0L;$PeakCombined=0L
+            $GpuProcessCounterAvailable=$false;$NextGpuSampleAt=[DateTime]::UtcNow
+            $OverMemoryLimit=$false
             try
             {
                 while(-not$Child.HasExited)
                 {
                     $Child.Refresh();if($Child.HasExited){break}
-                    $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+                    $Child.PriorityClass=$RequestedPriority;$Child.ProcessorAffinity=$AffinityPointer
                     $PeakPrivate=[Math]::Max($PeakPrivate,$Child.PrivateMemorySize64)
                     if($PeakPrivate-ge$EightGB)
                     {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+                    if([DateTime]::UtcNow-ge$NextGpuSampleAt)
+                    {
+                        $NextGpuSampleAt=[DateTime]::UtcNow.AddMilliseconds(500)
+                        try
+                        {
+                            $GpuSamples=@((Get-Counter -Counter '\GPU Process Memory(*)\Local Usage' -MaxSamples 1 -ErrorAction Stop).CounterSamples |
+                                Where-Object {$_.InstanceName-like "pid_$($Child.Id)_*"})
+                            if($GpuSamples.Count-gt0)
+                            {
+                                $GpuProcessCounterAvailable=$true
+                                $GpuLocalBytes=[long](($GpuSamples|Measure-Object -Property CookedValue -Sum).Sum)
+                                $PeakProcessGpuLocal=[Math]::Max($PeakProcessGpuLocal,$GpuLocalBytes)
+                                $CombinedNow=$Child.PrivateMemorySize64+$GpuLocalBytes
+                                $PeakCombined=[Math]::Max($PeakCombined,$CombinedNow)
+                                if($CombinedNow-ge$EightGB)
+                                {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+                            }
+                        }
+                        catch{}
+                    }
                     Start-Sleep -Milliseconds 100
                 }
             }
@@ -74,16 +105,21 @@ foreach($View in $Views)
             $Child.WaitForExit();$Child.Refresh();$ChildExitCode=[int]$Child.ExitCode
             if($ChildExitCode-eq-1073741502-and$Attempt-eq0){continue}
             if($ChildExitCode-eq-1073741502){throw '0xC0000142 repeated after its single permitted retry.'}
-            if($OverMemoryLimit){throw "Capture viewer stopped at the 8 GiB private-memory ceiling: $Stem"}
+            if($OverMemoryLimit)
+            {throw "Capture viewer stopped at memory ceiling: peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) process_gpu_local_peak_gb=$([math]::Round($PeakProcessGpuLocal/1GB,2)) sampled_combined_peak_gb=$([math]::Round($PeakCombined/1GB,2)) counter_available=$GpuProcessCounterAvailable capture=$Stem"}
             $LogText=Get-Content -LiteralPath $StdoutPath -Raw
             if(-not$LogText.Contains('[InitialClipmapReady]')-or-not(Test-Path -LiteralPath $CapturePath -PathType Leaf))
             {throw "Full-residency capture missing: $Stem; exit=$ChildExitCode"}
             $Memory=[regex]::Match($LogText,'(?m)^\[GpuFarMemory\] dxgi_local_sampled_peak_mb=([0-9.]+)')
-            $CombinedGB=$PeakPrivate/1GB
+            if($GpuProcessCounterAvailable){$CombinedGB=$PeakCombined/1GB}
-            if($Memory.Success){$CombinedGB=($PeakPrivate+[double]$Memory.Groups[1].Value*1MB)/1GB}
+            else
-            if($CombinedGB-ge8.0){throw "Capture exceeded conservative private + DXGI-local 8 GiB ceiling: $Stem ($CombinedGB GiB)"}
+            {
+                $CombinedGB=$PeakPrivate/1GB
+                if($Memory.Success){$CombinedGB=($PeakPrivate+[double]$Memory.Groups[1].Value*1MB)/1GB}
+            }
+            if($CombinedGB-ge8.0){throw "Capture exceeded private + process-local-GPU 8 GiB ceiling: $Stem ($CombinedGB GiB)"}
             $Ready=[regex]::Match($LogText,'(?m)^\[InitialClipmapReady\] ([0-9.]+) s')
-            Write-Output "STEP15_CAPTURE_DONE mode=$Mode view=$($View.Name) ready_s=$($Ready.Groups[1].Value) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) conservative_private_plus_dxgi_gb=$([math]::Round($CombinedGB,2)) path=$CapturePath exit=$ChildExitCode"
+            Write-Output "STEP15_CAPTURE_DONE mode=$Mode view=$($View.Name) ready_s=$($Ready.Groups[1].Value) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) process_gpu_local_peak_gb=$([math]::Round($PeakProcessGpuLocal/1GB,2)) peak_process_combined_gb=$([math]::Round($CombinedGB,2)) gpu_process_counter=$GpuProcessCounterAvailable path=$CapturePath exit=$ChildExitCode"
             if($ChildExitCode-ne0){throw "Capture viewer returned $ChildExitCode`: $Stem"}
             break
         }

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15CpuReference.ps1

@@ -1,24 +1,41 @@
 [CmdletBinding()]
-param([string]$ExecutableName='VoxelForgeGpuFieldStep15Timeline.exe',[string]$LogPrefix='step15_cpu_reference')
+param(
+    [string]$ExecutableName='VoxelForgeGpuFieldStep15Timeline.exe',
+    [string]$LogPrefix='step15_cpu_reference',
+    [string]$AffinityMask='0xFF00',
+    [switch]$NormalPriority,
+    [switch]$Include16Workers
+)
 
 $ErrorActionPreference='Stop'
 $HostProcess=Get-Process -Id $PID
-$HostProcess.PriorityClass='BelowNormal'
+$AffinityText=$AffinityMask.Trim()
-$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+if($AffinityText-match'^0[xX]([0-9a-fA-F]+)$'){$AffinityValue=[Convert]::ToUInt64($Matches[1],16)}
+elseif($AffinityText-match'^\d+$'){$AffinityValue=[Convert]::ToUInt64($AffinityText,10)}
+else{throw 'AffinityMask must be a non-zero decimal or hexadecimal processor mask.'}
+$AvailableAffinity=[uint64]$HostProcess.ProcessorAffinity.ToInt64()
+if($AffinityValue-eq0-or$AffinityValue-gt[uint64][long]::MaxValue-or
+    (($AffinityValue-band$AvailableAffinity)-ne$AffinityValue))
+{throw 'AffinityMask must fit the available logical processors.'}
+$AffinityPointer=[IntPtr][long]$AffinityValue
+$PriorityClass=if($NormalPriority){'Normal'}else{'BelowNormal'}
+$HostProcess.PriorityClass=$PriorityClass
+$HostProcess.ProcessorAffinity=$AffinityPointer
 $StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
 $ExecutablePath=Join-Path $StandaloneRoot ('bin\'+$ExecutableName)
 if([IO.Path]::GetFileName($ExecutableName)-ne$ExecutableName-or[IO.Path]::GetExtension($ExecutableName)-ne'.exe')
 {throw 'ExecutableName must be an executable filename.'}
 if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "Viewer not found: $ExecutablePath"}
 $Reference=$null
-foreach($Workers in @(1,8))
+$WorkerCounts=if($Include16Workers){@(1,8,16)}else{@(1,8)}
+foreach($Workers in $WorkerCounts)
 {
     $Name="${LogPrefix}_w$Workers"
     $StdoutPath=Join-Path $PSScriptRoot ($Name+'.stdout.log')
     $StderrPath=Join-Path $PSScriptRoot ($Name+'.stderr.log')
     $Arguments=@('--stream-benchmark','--shape-world','--force-generate','--stream-hash',
         '--clip-radius=3',"--workers=$Workers")
-    Write-Output "STEP15_CPU_REFERENCE workers=$Workers affinity=0xFF00 priority=BelowNormal command=`"$ExecutablePath`" $($Arguments -join ' ')"
+    Write-Output "STEP15_CPU_REFERENCE workers=$Workers affinity=$AffinityText priority=$PriorityClass command=`"$ExecutablePath`" $($Arguments -join ' ')"
     for($Attempt=0;$Attempt-lt2;$Attempt++)
     {
         if($Attempt-gt0)
@@ -32,7 +49,7 @@ foreach($Workers in @(1,8))
             while(-not$Child.HasExited)
             {
                 $Child.Refresh();if($Child.HasExited){break}
-                $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+                $Child.PriorityClass=$PriorityClass;$Child.ProcessorAffinity=$AffinityPointer
                 $PeakPrivate=[Math]::Max($PeakPrivate,$Child.PrivateMemorySize64)
                 if($Child.PrivateMemorySize64-ge(8L*1GB))
                 {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
@@ -55,7 +72,7 @@ foreach($Workers in @(1,8))
         $Hashes=@($Stream.Groups[1].Value.ToLowerInvariant(),$Stream.Groups[2].Value.ToLowerInvariant(),
             $Stream.Groups[3].Value.ToLowerInvariant(),$Stream.Groups[4].Value.ToLowerInvariant())
         if($null-eq$Reference){$Reference=$Hashes}
-        elseif(($Reference-join'|')-ne($Hashes-join'|')){throw 'CPU hashes differ between W1 and W8.'}
+        elseif(($Reference-join'|')-ne($Hashes-join'|')){throw "CPU hashes differ between requested worker counts through W$Workers."}
         Write-Output "STEP15_CPU_REFERENCE_RESULT workers=$Workers hashes=$($Hashes -join ',') walkability=5642/5642 peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) exit=$ChildExitCode"
         if($ChildExitCode-ne0){throw "CPU reference returned $ChildExitCode`: W$Workers"}
         break

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep15GpuSurvival.ps1

@@ -3,12 +3,22 @@ param([string]$ExecutableName='VoxelForgeGpuFieldStep15Timeline.exe',[string]$Lo
     [ValidateRange(10,600)][int]$NoProgressTimeoutSeconds=60,
     [ValidateSet('auto','cold','cached')][string]$WarmupPolicy='auto',
     [switch]$NoWarmupCacheWrite,
-    [switch]$SkipCpuDetailBand)
+    [string]$AffinityMask='0xFF00',
+    [switch]$NormalPriority)
 
 $ErrorActionPreference='Stop'
 $HostProcess=Get-Process -Id $PID
-$HostProcess.PriorityClass='BelowNormal'
+$AffinityText=$AffinityMask.Trim()
-$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+if($AffinityText -match '^0[xX]([0-9a-fA-F]+)$'){$AffinityValue=[Convert]::ToUInt64($Matches[1],16)}
+elseif($AffinityText -match '^\d+$'){$AffinityValue=[Convert]::ToUInt64($AffinityText,10)}
+else{throw 'AffinityMask must be a non-zero decimal or hexadecimal processor mask.'}
+$AvailableAffinity=[uint64]$HostProcess.ProcessorAffinity.ToInt64()
+if($AffinityValue-eq0-or$AffinityValue-gt[uint64][long]::MaxValue-or(($AffinityValue-band$AvailableAffinity)-ne$AffinityValue))
+{throw 'AffinityMask must fit the available logical processors.'}
+$AffinityPointer=[IntPtr][long]$AffinityValue
+$RequestedPriority=if($NormalPriority){'Normal'}else{'BelowNormal'}
+$HostProcess.PriorityClass=$RequestedPriority
+$HostProcess.ProcessorAffinity=$AffinityPointer
 $StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
 $ExecutablePath=Join-Path $StandaloneRoot ('bin\'+$ExecutableName)
 if([IO.Path]::GetFileName($ExecutableName)-ne$ExecutableName-or[IO.Path]::GetExtension($ExecutableName)-ne'.exe')
@@ -22,8 +32,7 @@ $EightGB=8L*1024L*1024L*1024L
 if($WarmupPolicy-eq'cold'){$Arguments+='--gpu-far-force-cold-start'}
 elseif($WarmupPolicy-eq'cached'){$Arguments+='--gpu-far-force-cached-start'}
 if($NoWarmupCacheWrite){$Arguments+='--gpu-far-no-warmup-cache-write'}
-if($SkipCpuDetailBand){$Arguments+='--gpu-far-skip-cpu-detail-band'}
+Write-Output "STEP15_SURVIVAL_RUN affinity=0x$('{0:X}' -f $AffinityValue) priority=$RequestedPriority command=`"$ExecutablePath`" $($Arguments -join ' ')"
-Write-Output "STEP15_SURVIVAL_RUN skip_cpu_detail_band=$([bool]$SkipCpuDetailBand) command=`"$ExecutablePath`" $($Arguments -join ' ')"
 Write-Output "STEP15_SURVIVAL_POLICY warmup=$WarmupPolicy no_warmup_cache_write=$([bool]$NoWarmupCacheWrite) no_progress_timeout_seconds=$NoProgressTimeoutSeconds"
 for($Attempt=0;$Attempt-lt2;$Attempt++)
 {
@@ -32,17 +41,39 @@ for($Attempt=0;$Attempt-lt2;$Attempt++)
     Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
     $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
         -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
-    $PeakPrivate=0L;$OverMemoryLimit=$false
+    $PeakPrivate=0L;$PeakProcessGpuLocal=0L;$PeakCombined=0L
+    $GpuProcessCounterAvailable=$false;$NextGpuSampleAt=[DateTime]::UtcNow
+    $OverMemoryLimit=$false
     $NoProgressTimedOut=$false;$LastProgressAt=[DateTime]::UtcNow;$LastOutputWrite=$null
     try
     {
         while(-not$Child.HasExited)
         {
             $Child.Refresh();if($Child.HasExited){break}
-            $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+            $Child.PriorityClass=$RequestedPriority;$Child.ProcessorAffinity=$AffinityPointer
             $PeakPrivate=[Math]::Max($PeakPrivate,$Child.PrivateMemorySize64)
             if($Child.PrivateMemorySize64-ge$EightGB)
             {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+            if([DateTime]::UtcNow-ge$NextGpuSampleAt)
+            {
+                $NextGpuSampleAt=[DateTime]::UtcNow.AddMilliseconds(500)
+                try
+                {
+                    $GpuSamples=@((Get-Counter -Counter '\GPU Process Memory(*)\Local Usage' -MaxSamples 1 -ErrorAction Stop).CounterSamples |
+                        Where-Object {$_.InstanceName-like "pid_$($Child.Id)_*"})
+                    if($GpuSamples.Count-gt0)
+                    {
+                        $GpuProcessCounterAvailable=$true
+                        $GpuLocalBytes=[long](($GpuSamples|Measure-Object -Property CookedValue -Sum).Sum)
+                        $PeakProcessGpuLocal=[Math]::Max($PeakProcessGpuLocal,$GpuLocalBytes)
+                        $CombinedNow=$Child.PrivateMemorySize64+$GpuLocalBytes
+                        $PeakCombined=[Math]::Max($PeakCombined,$CombinedNow)
+                        if($CombinedNow-ge$EightGB)
+                        {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+                    }
+                }
+                catch{}
+            }
             if(Test-Path -LiteralPath $StdoutPath -PathType Leaf)
             {
                 $OutputWrite=(Get-Item -LiteralPath $StdoutPath).LastWriteTimeUtc
@@ -61,7 +92,7 @@ for($Attempt=0;$Attempt-lt2;$Attempt++)
     $Child.WaitForExit();$Child.Refresh();$ChildExitCode=[int]$Child.ExitCode
     if($ChildExitCode-eq-1073741502-and$Attempt-eq0){continue}
     if($ChildExitCode-eq-1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
-    if($OverMemoryLimit){throw 'GPU survival audit stopped at the private-memory ceiling.'}
+    if($OverMemoryLimit){throw "GPU survival stopped at memory ceiling: peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) process_gpu_local_peak_gb=$([math]::Round($PeakProcessGpuLocal/1GB,2)) sampled_combined_peak_gb=$([math]::Round($PeakCombined/1GB,2)) counter_available=$GpuProcessCounterAvailable"}
     if($NoProgressTimedOut)
     {
         Write-Output "STEP15_SURVIVAL_WATCHDOG state=no_stdout_progress timeout_seconds=$NoProgressTimeoutSeconds exit=$ChildExitCode stdout=$StdoutPath stderr=$StderrPath"
@@ -75,9 +106,10 @@ for($Attempt=0;$Attempt-lt2;$Attempt++)
     $Memory=[regex]::Match($LogText,'(?m)^\[GpuFarMemory\] dxgi_local_sampled_peak_mb=([0-9.]+)')
     if(-not$Audit.Success){throw 'GPU survival audit result marker missing.'}
     $DxgiMB=if($Memory.Success){[double]$Memory.Groups[1].Value}else{0.0}
-    $CombinedGB=($PeakPrivate+$DxgiMB*1MB)/1GB
+    if($GpuProcessCounterAvailable){$CombinedGB=$PeakCombined/1GB}
-    if($CombinedGB-ge8.0){throw "GPU survival exceeded conservative private + DXGI-local 8 GiB ceiling: $CombinedGB GiB"}
+    else{$CombinedGB=($PeakPrivate+$DxgiMB*1MB)/1GB}
-    Write-Output "STEP15_SURVIVAL_RESULT audit=$($Audit.Groups[1].Value) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) dxgi_local_peak_mb=$DxgiMB conservative_private_plus_dxgi_gb=$([math]::Round($CombinedGB,2)) exit=$ChildExitCode stdout=$StdoutPath stderr=$StderrPath"
+    if($CombinedGB-ge8.0){throw "GPU survival exceeded private + process-local-GPU 8 GiB ceiling: $CombinedGB GiB"}
+    Write-Output "STEP15_SURVIVAL_RESULT audit=$($Audit.Groups[1].Value) peak_private_gb=$([math]::Round($PeakPrivate/1GB,2)) process_gpu_local_peak_mb=$([math]::Round($PeakProcessGpuLocal/1MB,1)) dxgi_local_adapter_peak_mb=$DxgiMB peak_process_combined_gb=$([math]::Round($CombinedGB,2)) gpu_process_counter=$GpuProcessCounterAvailable exit=$ChildExitCode stdout=$StdoutPath stderr=$StderrPath"
     if($Audit.Groups[1].Value-ne'PASS'){throw 'GPU survival audit failed.'}
     if($ChildExitCode-ne0){throw "Viewer returned $ChildExitCode"}
     break

Standalone/Research/OwnerShapeWorld/gpu_field_port/STEP28_REPORT.md

@@ -0,0 +1,81 @@
+# Step 28 report: GPU near skirts and compact panel output
+
+Date: 2026-10-08
+
+All implementation, build, validation, capture, and report artifacts are under `Standalone\`. No cache or driver state was manually changed. Capture files were compared as numeric RGB arrays and were not opened in an image viewer. No commit, push, or stash was made.
+
+## Shader build
+
+The viewer was built with the complete FXC `/O3` shader pass using:
+
+```powershell
+BuildGpuFieldViewer.ps1 -OutputName VoxelForgeGpuFarStep28.exe -AffinityMask 0xFFFF -NormalPriority
+```
+
+All 20 production compute entries compiled, including `ClassifyFarSkirtsCS`, `ScanFarSkirtsCS`, and `EmitFarSkirtsCS`. The skirt kernels that had previously run for more than ten hours completed in the ordinary full build. The build also printed `CORE_LINK_SAFETY_PASS` and produced `Standalone\bin\VoxelForgeGpuFarStep28.exe`. The agreement runner built successfully as `Standalone\bin\VoxelForgeGpuFarAgreementStep28.exe`.
+
+Near skirts now use the same near-resolution refinement decisions as the CPU path. The GPU pipeline first prepares and evaluates the near child probes, then the skirt classify/scan/emit work consumes those results. This keeps the costly near field evaluator out of the skirt kernels and lets FXC compile the full shader at `/O3`.
+
+The viewer allocates panel/skirt output ranges from prepared input data: each range reserves its exact connector vertex count plus a guarded skirt capacity when the tile has visible skirt faces. The pool had 11,729,896 vertices for 2,074 full-residency candidates, compared with 67,960,832 vertices from the former fixed 32,768-vertex allocation per candidate (82.7% fewer reserved vertices). Normal output status checks still fall back to CPU if a tile exceeds its range.
+
+## CPU reference and agreement
+
+W1, W8, and W16 all matched the four fixed hashes, with walkability `5642/5642`:
+
+| Hash | Reference |
+|---|---|
+| Density | `0xdb2991c73215af4c` |
+| Mesh | `0x8236c5edd7fc8e4f` |
+| Payload | `0x4c48fa75a543d58e` |
+| Topology | `0xbea22e76b34a4dfc` |
+
+Whole-near agreement covered 960 tiles and 640 visible skirt audits. CPU, hardware, and WARP had zero sign differences, topology differences, protected exception differences, skirt overflow, or skirt audit failures. One skirt tile, `(-3,-10,0,lod1)`, had 16 triangles with different one-centimetre quantized keys. Its maximum vertex displacement was 0.00109183 cm; its 0.15 cm triangle topology comparison had zero differences. The overall runner reported `PASS`.
+
+Whole-far agreement covered 1,114 tile records, 127,401 protected exception cells, and 353 thin-feature cells. Hardware and WARP each had zero sign, topology, protected, or thin-feature differences. The targeted thin-feature selector chose `(-1,0,-1,lod3)`; it audited 35 protected exception cells and 18 thin-feature cells, with zero hardware or WARP differences. All agreement runs reported `far_field_topology_protected_status=PASS`.
+
+## Full-residency survival and memory
+
+Both auto and forced-cold GPU survival runs passed `[GpuFarSurvivalAudit]` and `[FarSurvivalAudit]`. Each checked 2,074 tiles; surface overflow, panel overflow, and count-readback failures were all zero. The maximum required surface output was 18,153 of 19,456 vertices at `(2,-9,-4,lod1)`, slot 714. The far audit found 7,200/7,200 boundary surfaces and 10,976/10,976 good connectors.
+
+| Route | Peak private | Peak viewer GPU-local | Sampled combined |
+|---|---:|---:|---:|
+| Auto | 4.17 GiB | 3,135.3 MiB | 7.00 GiB |
+| Forced cold | 4.08 GiB | 3,128.3 MiB | 7.02 GiB |
+
+The process-specific `GPU Process Memory` counter was available for both runs. The 8 GiB per-process ceiling held in both cases.
+
+## Ten-view capture comparison
+
+All ten CPU/GPU pairs reached `[InitialClipmapReady]`. Numeric comparison found zero CPU-rock-to-GPU-background pixels and zero GPU-rock-to-CPU-background pixels in every view. The report contains RGB MAE/RMSE and changed-pixel metrics: [step28_final_capture_comparison.md](step28_final_capture_comparison.md).
+
+The capture runner originally rejected the first GPU image because it added whole-adapter DXGI-local usage to viewer private memory. I updated it to sample the viewer PID's GPU-local counter, consistent with the survival runner, then reran all ten pairs successfully. Peak per-process checks passed; this removes the adapter-wide overcount from the capture memory gate.
+
+## CPU detail-band skip retired
+
+The Step 27 `-SkipCpuDetailBand` option was removed from the viewer, agreement, survival, and capture runners. In the attempted whole-near skip run, 154 of 640 skirt audits failed, with vertex/topology errors as large as 4.69 cm. The option also added about 174 ms to interval/sign planning in the owner's race. That planning cost outweighed its detail-band saving, and the option did not meet the exact skirt behavior requirement.
+
+## Timing
+
+The comparisons used eight workers. The 0xFF00 runs used BelowNormal priority; the 0xFFFF runs used Normal priority.
+
+| Mask and route | Initial clipmap ready | GPU tiles submitted | Interpretation |
+|---|---:|---:|---|
+| 0xFF00 CPU | 1.348 s | 0 | CPU reference route |
+| 0xFF00 GPU auto | 1.370 s | 0 | `cache_marker=miss`; CPU fallback, not a cached-GPU timing |
+| 0xFF00 GPU warmup before workers, run 1 | 3.508 s | 894 near, 0 far | Cold warm-up; remaining candidates went through CPU |
+| 0xFF00 GPU warmup before workers, run 2 | 2.916 s | 960 near, 710 far | Cold warm-up; all near tiles submitted |
+| 0xFFFF CPU | 1.300 s | 0 | CPU reference route |
+| 0xFFFF GPU auto | 1.303 s | 0 | `cache_marker=miss`; CPU fallback, not a cached-GPU timing |
+| 0xFFFF GPU warmup before workers | 2.922 s | 960 near, 1,106 far | Cold warm-up; 160 far tiles used CPU fallback |
+
+The ten-view 0xFFFF capture run separately reached readiness in 2.720 s and submitted 960 near and 1,114 far tiles. Auto marker-miss times closely matched CPU, but neither auto run submitted a GPU tile. The cached-GPU timing target is therefore still unverified in this workspace. No cache write was enabled for these measurements.
+
+## Logs and captures
+
+- Agreement: `step28_compact_whole_near.*`, `step28_compact_whole_far.*`, `step28_compact_thin_feature.*`
+- Survival: `step28_compact_survival_auto.*`, `step28_compact_survival_cold.*`
+- Ten capture pairs: `step28_final_{cpu,gpu}_*.png` and [the numeric comparison](step28_final_capture_comparison.md)
+- Timings: `step28_timing_*.stdout.log` and matching stderr logs
+- CPU reference: `step28_final_cpu_reference_w{1,8,16}.stdout.log`
+
+See [RUNNERS.md](RUNNERS.md) for the current commands and [LUNA_DIALOGUE.md](LUNA_DIALOGUE.md) for the step 28 dialogue entry.

Standalone/Research/OwnerShapeWorld/gpu_field_port/step28_final_capture_comparison.md

@@ -0,0 +1,18 @@
+# Step28_Final full-residency capture comparison
+
+All ten CPU/GPU pairs were captured after `[InitialClipmapReady]`. The files were compared numerically as RGB arrays and were not opened in an image viewer.
+
+`CPU rock to GPU background` isolates the prior missing-coverage signature. Captures were loaded as RGB arrays only; no full-size image was opened in a viewer.
+
+| View | Size | RGB MAE | RGB RMSE | Changed pixels | Pixels max channel >32 | Max channel | CPU rock to GPU background | GPU rock to CPU background | Missing-pixel bounds |
+|---|---:|---:|---:|---:|---:|---:|---:|---:|---|
+| chamber | 1280x800 | 0.6075 | 1.8638 | 42.6368% | 0.0183% | 123 | 0 | 0 | none |
+| tunnel | 1280x800 | 0.7641 | 3.4370 | 41.7169% | 0.2778% | 130 | 0 | 0 | none |
+| passage | 1280x800 | 1.6092 | 3.6878 | 50.4982% | 0.1733% | 52 | 0 | 0 | none |
+| far | 1280x800 | 0.7910 | 2.1578 | 44.8325% | 0.0198% | 128 | 0 | 0 | none |
+| passage_floor_closeup | 1280x800 | 0.9362 | 2.4902 | 43.1415% | 0.0513% | 100 | 0 | 0 | none |
+| tunnel_bend | 1280x800 | 0.8234 | 2.1697 | 41.4529% | 0.0708% | 55 | 0 | 0 | none |
+| remote_yaw0 | 1280x800 | 0.0491 | 0.9228 | 3.9361% | 0.0274% | 92 | 0 | 0 | none |
+| remote_yaw90 | 1280x800 | 0.0310 | 0.6503 | 2.8547% | 0.0073% | 97 | 0 | 0 | none |
+| remote_yaw180 | 1280x800 | 0.0582 | 0.7545 | 5.2343% | 0.0090% | 106 | 0 | 0 | none |
+| remote_yaw270 | 1280x800 | 0.0830 | 1.1768 | 6.2815% | 0.0323% | 85 | 0 | 0 | none |

Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl

@@ -32,6 +32,11 @@ cbuffer TileConstants : register(b0)
     uint3 TileConstantsPadding;
 };
 
+uint FarSkirtNearProbeSurfaceCapacity()
+{
+    return TileConstantsPadding.x;
+}
+
 cbuffer FarFieldConstants : register(b1)
 {
     float FarStrataLo;
@@ -62,7 +67,7 @@ cbuffer FarConnectorConstants : register(b2)
     uint FarShapeCount;
     uint FarSkirtFaceMask;
     float FarSkirtDepthVoxels;
-    uint FarConnectorPadding;
+    uint FarConnectorOutputBaseVertex;
 };
 
 static const int3 CornerOffset[8] = {
@@ -448,6 +453,7 @@ NearFieldBaseData NearFieldBaseAt(float3 position,out uint distanceEvaluations)
     const FarFieldSpatialBin range=FarFieldFindBin(position);
     float cave=100000.0f,overlay=-100000.0f;
     float caveFloorInfluence=0.0f,cavePassageInfluence=0.0f;
+    bool criticalAir=false;
     distanceEvaluations=0;
     [loop]
     for(uint item=0;item<range.count;++item)
@@ -455,6 +461,22 @@ NearFieldBaseData NearFieldBaseAt(float3 position,out uint distanceEvaluations)
         const VF_FieldPrimitive shape=FarShapes[FarSpatialBinShapeIds[range.offset+item]];
         if((shape.flags&VF_FIELD_PRIMITIVE_CARVE)!=0)
         {
+            if((shape.flags&VF_FIELD_PRIMITIVE_CRITICAL)!=0
+                &&all(position>=float3(shape.minVoxel[0],shape.minVoxel[1],shape.minVoxel[2]))
+                &&all(position<=float3(shape.maxVoxel[0],shape.maxVoxel[1],shape.maxVoxel[2])))
+            {
+                const int3 decidedPosition=int3(position);
+                float decidedDistance=VF_FieldPrimitiveDistance(shape,
+                    (float)decidedPosition.x,(float)decidedPosition.y,(float)decidedPosition.z);
+                if((shape.flags&VF_FIELD_PRIMITIVE_HAS_FLOOR)!=0)
+                {
+                    const float decidedFloor=VF_FieldPrimitiveFloorAt(shape,
+                        (float)decidedPosition.x,(float)decidedPosition.y);
+                    decidedDistance=VF_FieldSmoothMax(decidedDistance,
+                        decidedFloor-(float)decidedPosition.z,shape.floorBlendRadius);
+                }
+                criticalAir=criticalAir||decidedDistance<=0.0f;
+            }
             if(position.x<shape.minVoxel[0]-6||position.x>shape.maxVoxel[0]+6
                 ||position.y<shape.minVoxel[1]-6||position.y>shape.maxVoxel[1]+6
                 ||position.z<shape.minVoxel[2]-6||position.z>shape.maxVoxel[2]+6)continue;
@@ -495,7 +517,7 @@ NearFieldBaseData NearFieldBaseAt(float3 position,out uint distanceEvaluations)
     const float rockDistance=VF_FieldSmoothMin(base,carveOrOverlay,FarStrataBlend);
     NearFieldBaseData result;
     result.geometry=float4(rockDistance,surfaceSdf,caveInfluence,floorInfluence);
-    result.influences=float4(passageInfluence,0,0,0);
+    result.influences=float4(passageInfluence,criticalAir?1.0f:0.0f,0,0);
     return result;
 }
 
@@ -567,10 +589,12 @@ float NearFieldDetailRockDistance(float3 position,NearFieldBaseData baseData)
     return rockDistance;
 }
 
+[noinline]
 float NearFieldDensityAt(float3 position,uint clampNearZero,out uint distanceEvaluations)
 {
     const NearFieldBaseData baseData=NearFieldBaseAt(position,distanceEvaluations);
-    const float rockDistance=NearFieldDetailRockDistance(position,baseData);
+    float rockDistance=NearFieldDetailRockDistance(position,baseData);
+    if(baseData.influences.y!=0.0f)rockDistance=min(rockDistance,-1.0f/16.0f);
     float density=-rockDistance;
     if(clampNearZero!=0&&abs(density)<0.05f)
         density=rockDistance>=0.0f?-0.05f:0.05f;
@@ -757,6 +781,48 @@ struct FarSkirtCachedSamples
 
 StructuredBuffer<FarSkirtCachedSamples> FarSkirtCachedProbeInput : register(t21);
 RWStructuredBuffer<FarSkirtCachedSamples> FarSkirtCachedProbeOutput : register(u2);
+StructuredBuffer<float4> FarSkirtNearProbeInput : register(t22);
+ByteAddressBuffer FarSkirtNearProbeCounterInput : register(t23);
+RWStructuredBuffer<float4> FarSkirtNearProbeOutput : register(u3);
+RWByteAddressBuffer FarSkirtNearProbeCounter : register(u5);
+RWByteAddressBuffer FarSkirtNearProbeDispatchArgs : register(u6);
+
+static const uint FarSkirtNearRootProbeStride=11u;
+static const uint FarSkirtNearChildProbeFanStride=4u;
+static const uint FarSkirtNearChildProbeFanLimit=8u;
+static const uint FarSkirtNearRootLimit=6u;
+static const uint FarSkirtNearProbeStride=FarSkirtNearRootProbeStride
+    +FarSkirtNearChildProbeFanStride*FarSkirtNearChildProbeFanLimit;
+static const float FarSkirtUnusedNearProbeDensity=-3.402823e+38f;
+
+[noinline]
+float NearSkirtDensityAtVoxel(float3 position)
+{
+    uint ignoredEvaluations=0;
+    float density=NearFieldDensityAt(position,1,ignoredEvaluations);
+    const float3 latticePosition=(position-float3(OriginX,OriginY,OriginZ))/StepVoxels;
+    const int3 lattice=int3(round(latticePosition));
+    if(all(abs(latticePosition-float3(lattice))<=1.0e-4f)
+        &&all(lattice>=-1)&&all(lattice<=int(CellsPerAxis)+1))
+    {
+        const uint sample=SampleIndex(lattice);
+        const uint packedProxy=FarNearPackedProxy[sample];
+        if(packedProxy!=0xffffffffu)
+        {
+            const bool packedSignOverride=(packedProxy&0x40000000u)!=0u;
+            const bool packedRockSign=(packedProxy&0x80000000u)!=0u;
+            const uint magnitudeQ=packedProxy&0x3fffffffu;
+            const float magnitude=magnitudeQ==0u?4.125f:(float)magnitudeQ*(1.0f/256.0f);
+            const bool negative=packedSignOverride?packedRockSign:(density<=0.0f);
+            density=negative?-magnitude:magnitude;
+        }
+        const uint protectedAir=(FarProtectedAir[sample>>5]>>(sample&31u))&1u;
+        const uint protectedSolid=(FarProtectedSolid[sample>>5]>>(sample&31u))&1u;
+        if(protectedAir!=0u)density=max(abs(density),1.0f/256.0f);
+        else if(protectedSolid!=0u)density=-max(abs(density),1.0f/256.0f);
+    }
+    return density;
+}
 
 float FarSkirtDensityAtVoxel(float3 position)
 {
@@ -783,8 +849,8 @@ FarSkirtVertex FarSkirtInterpolate(FarSkirtVertex a,FarSkirtVertex b,float t)
     FarSkirtVertex result;
     result.position=lerp(a.position,b.position,t);
     result.normal=normalize(lerp(a.normal,b.normal,t));
-    // Clipping only uses the interpolated position and normal after insertion.
+    // Clipping uses the interpolated density only to find this root. Later near
-    // Its density is never fed back into an edge decision or emitted vertex.
+    // refinement reads its exact probe from FarSkirtNearProbeInput.
     result.density=lerp(a.density,b.density,t);
     return result;
 }
@@ -812,8 +878,221 @@ void FarSkirtClipTriangle(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex c,
     }
 }
 
+void FarSkirtStoreNearRootProbePositions(FarSkirtVertex a,FarSkirtVertex b,
+    FarSkirtVertex c,uint probeBase)
+{
+    const FarSkirtVertex midpoint01=FarSkirtInterpolate(a,b,0.5f);
+    const FarSkirtVertex midpoint12=FarSkirtInterpolate(b,c,0.5f);
+    const FarSkirtVertex midpoint20=FarSkirtInterpolate(c,a,0.5f);
+    FarSkirtNearProbeOutput[probeBase]=float4(a.position/VoxelSizeCm,0.0f);
+    FarSkirtNearProbeOutput[probeBase+1]=float4(b.position/VoxelSizeCm,0.0f);
+    FarSkirtNearProbeOutput[probeBase+2]=float4(c.position/VoxelSizeCm,0.0f);
+    FarSkirtNearProbeOutput[probeBase+3]=float4(
+        (a.position+b.position+c.position)/(3.0f*VoxelSizeCm),0.0f);
+    FarSkirtNearProbeOutput[probeBase+4]=float4(midpoint01.position/VoxelSizeCm,0.0f);
+    FarSkirtNearProbeOutput[probeBase+5]=float4(midpoint12.position/VoxelSizeCm,0.0f);
+    FarSkirtNearProbeOutput[probeBase+6]=float4(midpoint20.position/VoxelSizeCm,0.0f);
+    FarSkirtNearProbeOutput[probeBase+7]=float4(
+        (a.position+midpoint01.position+midpoint20.position)/(3.0f*VoxelSizeCm),0.0f);
+    FarSkirtNearProbeOutput[probeBase+8]=float4(
+        (midpoint01.position+b.position+midpoint12.position)/(3.0f*VoxelSizeCm),0.0f);
+    FarSkirtNearProbeOutput[probeBase+9]=float4(
+        (midpoint20.position+midpoint12.position+c.position)/(3.0f*VoxelSizeCm),0.0f);
+    FarSkirtNearProbeOutput[probeBase+10]=float4(
+        (midpoint01.position+midpoint12.position+midpoint20.position)/(3.0f*VoxelSizeCm),0.0f);
+    [unroll]
+    for(uint childProbe=FarSkirtNearRootProbeStride;
+        childProbe<FarSkirtNearProbeStride;++childProbe)
+        FarSkirtNearProbeOutput[probeBase+childProbe]=float4(0.0f,0.0f,0.0f,
+            FarSkirtUnusedNearProbeDensity);
+}
+
+void FarSkirtPrepareNearSourceTriangle(FarSkirtVertex a,FarSkirtVertex b,
+    FarSkirtVertex c,uint probeBase,inout uint rootIndex)
+{
+    FarSkirtVertex polygon[5];uint polygonCount=0;
+    FarSkirtClipTriangle(a,b,c,polygon,polygonCount);
+    if(polygonCount<3)return;
+    [loop]
+    for(uint fan=1;fan+1<polygonCount;++fan)
+    {
+        if(rootIndex<FarSkirtNearRootLimit)
+            FarSkirtStoreNearRootProbePositions(polygon[0],polygon[fan],polygon[fan+1],
+                probeBase+rootIndex*FarSkirtNearProbeStride);
+        ++rootIndex;
+    }
+}
+
+FarSkirtVertex FarSkirtNearRootVertex(FarSkirtVertex vertex,float exactDensity,
+    FarSkirtVertex sourceA,FarSkirtVertex sourceB)
+{
+    vertex.density=all(vertex.position==sourceA.position)||all(vertex.position==sourceB.position)
+        ?0.0f:exactDensity;
+    return vertex;
+}
+
+void FarSkirtNearBuildChildTriangleOutput(FarSkirtVertex a,FarSkirtVertex b,
+    FarSkirtVertex c,uint probeBase,uint rootIndex,FarSkirtVertex sourceA,
+    FarSkirtVertex sourceB,uint childIndex,out FarSkirtVertex childA,
+    out FarSkirtVertex childB,out FarSkirtVertex childC)
+{
+    const uint base=probeBase+rootIndex*FarSkirtNearProbeStride;
+    a=FarSkirtNearRootVertex(a,FarSkirtNearProbeOutput[base].w,sourceA,sourceB);
+    b=FarSkirtNearRootVertex(b,FarSkirtNearProbeOutput[base+1].w,sourceA,sourceB);
+    c=FarSkirtNearRootVertex(c,FarSkirtNearProbeOutput[base+2].w,sourceA,sourceB);
+    FarSkirtVertex midpoint01=FarSkirtInterpolate(a,b,0.5f);
+    FarSkirtVertex midpoint12=FarSkirtInterpolate(b,c,0.5f);
+    FarSkirtVertex midpoint20=FarSkirtInterpolate(c,a,0.5f);
+    midpoint01.density=FarSkirtNearProbeOutput[base+4].w;
+    midpoint12.density=FarSkirtNearProbeOutput[base+5].w;
+    midpoint20.density=FarSkirtNearProbeOutput[base+6].w;
+    if(childIndex==0){childA=a;childB=midpoint01;childC=midpoint20;}
+    else if(childIndex==1){childA=midpoint01;childB=b;childC=midpoint12;}
+    else if(childIndex==2){childA=midpoint20;childB=midpoint12;childC=c;}
+    else {childA=midpoint01;childB=midpoint12;childC=midpoint20;}
+}
+
+void FarSkirtStoreNearChildFanProbePositions(FarSkirtVertex a,FarSkirtVertex b,
+    FarSkirtVertex c,uint probeBase,uint rootIndex,uint childFanIndex)
+{
+    const uint base=probeBase+rootIndex*FarSkirtNearProbeStride
+        +FarSkirtNearRootProbeStride+childFanIndex*FarSkirtNearChildProbeFanStride;
+    FarSkirtNearProbeOutput[base]=float4(a.position/VoxelSizeCm,0.0f);
+    FarSkirtNearProbeOutput[base+1]=float4(b.position/VoxelSizeCm,0.0f);
+    FarSkirtNearProbeOutput[base+2]=float4(c.position/VoxelSizeCm,0.0f);
+    FarSkirtNearProbeOutput[base+3]=float4(
+        (a.position+b.position+c.position)/(3.0f*VoxelSizeCm),0.0f);
+}
+
+float FarSkirtNearRootMaximumOutput(uint probeBase,uint rootIndex)
+{
+    const uint base=probeBase+rootIndex*FarSkirtNearProbeStride;
+    float maximumDensity=FarSkirtNearProbeOutput[base].w;
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeOutput[base+1].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeOutput[base+2].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeOutput[base+3].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeOutput[base+4].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeOutput[base+5].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeOutput[base+6].w);
+    return maximumDensity;
+}
+
+void FarSkirtPrepareNearChildRoot(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex c,
+    uint probeBase,uint rootIndex,FarSkirtVertex sourceA,FarSkirtVertex sourceB)
+{
+    if(FarSkirtNearRootMaximumOutput(probeBase,rootIndex)<=0.25f)return;
+    uint childFanIndex=0;
+    [unroll]
+    for(uint child=0;child<4;++child)
+    {
+        FarSkirtVertex childA,childB,childC;
+        FarSkirtNearBuildChildTriangleOutput(a,b,c,probeBase,rootIndex,sourceA,sourceB,
+            child,childA,childB,childC);
+        FarSkirtVertex polygon[5];uint polygonCount=0;
+        FarSkirtClipTriangle(childA,childB,childC,polygon,polygonCount);
+        if(polygonCount<3)continue;
+        [loop]
+        for(uint fan=1;fan+1<polygonCount;++fan)
+            FarSkirtStoreNearChildFanProbePositions(polygon[0],polygon[fan],polygon[fan+1],
+                probeBase,rootIndex,childFanIndex++);
+    }
+}
+
+void FarSkirtPrepareNearChildSourceTriangle(FarSkirtVertex a,FarSkirtVertex b,
+    FarSkirtVertex c,FarSkirtVertex sourceA,FarSkirtVertex sourceB,uint probeBase,
+    inout uint rootIndex)
+{
+    FarSkirtVertex polygon[5];uint polygonCount=0;
+    FarSkirtClipTriangle(a,b,c,polygon,polygonCount);
+    if(polygonCount<3)return;
+    [loop]
+    for(uint fan=1;fan+1<polygonCount;++fan)
+    {
+        if(rootIndex<FarSkirtNearRootLimit)
+            FarSkirtPrepareNearChildRoot(polygon[0],polygon[fan],polygon[fan+1],
+                probeBase,rootIndex,sourceA,sourceB);
+        ++rootIndex;
+    }
+}
+
+uint FarSkirtCountNearSourceTriangleRoots(FarSkirtVertex a,FarSkirtVertex b,
+    FarSkirtVertex c)
+{
+    FarSkirtVertex polygon[5];uint polygonCount=0;
+    FarSkirtClipTriangle(a,b,c,polygon,polygonCount);
+    return polygonCount<3?0:polygonCount-2;
+}
+
+float FarSkirtNearRootMaximum(uint probeBase,uint rootIndex)
+{
+    const uint base=probeBase+rootIndex*FarSkirtNearProbeStride;
+    float maximumDensity=FarSkirtNearProbeInput[base].w;
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeInput[base+1].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeInput[base+2].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeInput[base+3].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeInput[base+4].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeInput[base+5].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeInput[base+6].w);
+    return maximumDensity;
+}
+
+void FarSkirtNearBuildChildTriangle(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex c,
+    uint probeBase,uint rootIndex,FarSkirtVertex sourceA,FarSkirtVertex sourceB,
+    uint childIndex,out FarSkirtVertex childA,out FarSkirtVertex childB,
+    out FarSkirtVertex childC)
+{
+    const uint base=probeBase+rootIndex*FarSkirtNearProbeStride;
+    a=FarSkirtNearRootVertex(a,FarSkirtNearProbeInput[base].w,sourceA,sourceB);
+    b=FarSkirtNearRootVertex(b,FarSkirtNearProbeInput[base+1].w,sourceA,sourceB);
+    c=FarSkirtNearRootVertex(c,FarSkirtNearProbeInput[base+2].w,sourceA,sourceB);
+    FarSkirtVertex midpoint01=FarSkirtInterpolate(a,b,0.5f);
+    FarSkirtVertex midpoint12=FarSkirtInterpolate(b,c,0.5f);
+    FarSkirtVertex midpoint20=FarSkirtInterpolate(c,a,0.5f);
+    midpoint01.density=FarSkirtNearProbeInput[base+4].w;
+    midpoint12.density=FarSkirtNearProbeInput[base+5].w;
+    midpoint20.density=FarSkirtNearProbeInput[base+6].w;
+    if(childIndex==0){childA=a;childB=midpoint01;childC=midpoint20;}
+    else if(childIndex==1){childA=midpoint01;childB=b;childC=midpoint12;}
+    else if(childIndex==2){childA=midpoint20;childB=midpoint12;childC=c;}
+    else {childA=midpoint01;childB=midpoint12;childC=midpoint20;}
+}
+
+uint FarSkirtNearChildFanOutputCount(uint probeBase,uint rootIndex,uint childFanIndex)
+{
+    const uint base=probeBase+rootIndex*FarSkirtNearProbeStride
+        +FarSkirtNearRootProbeStride+childFanIndex*FarSkirtNearChildProbeFanStride;
+    float maximumDensity=FarSkirtNearProbeInput[base].w;
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeInput[base+1].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeInput[base+2].w);
+    maximumDensity=max(maximumDensity,FarSkirtNearProbeInput[base+3].w);
+    return maximumDensity<=0.25f?6u:0u;
+}
+
+uint FarSkirtNearResolutionTriangleOutputCount(FarSkirtVertex a,FarSkirtVertex b,
+    FarSkirtVertex c,uint probeBase,uint rootIndex,FarSkirtVertex sourceA,
+    FarSkirtVertex sourceB)
+{
+    if(FarSkirtNearRootMaximum(probeBase,rootIndex)<=0.25f)return 6;
+    uint vertexCount=0,childFanIndex=0;
+    [unroll]
+    for(uint child=0;child<4;++child)
+    {
+        FarSkirtVertex childA,childB,childC;
+        FarSkirtNearBuildChildTriangle(a,b,c,probeBase,rootIndex,sourceA,sourceB,
+            child,childA,childB,childC);
+        FarSkirtVertex polygon[5];uint polygonCount=0;
+        FarSkirtClipTriangle(childA,childB,childC,polygon,polygonCount);
+        if(polygonCount<3)continue;
+        [loop]
+        for(uint fan=1;fan+1<polygonCount;++fan)
+            vertexCount+=FarSkirtNearChildFanOutputCount(probeBase,rootIndex,childFanIndex++);
+    }
+    return vertexCount;
+}
+
 uint FarSkirtSourceTriangleOutputCount(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex c,
-    inout FarSkirtCachedSamples cached)
+    inout FarSkirtCachedSamples cached,uint nearProbeBase,inout uint nearRootIndex,
+    FarSkirtVertex sourceA,FarSkirtVertex sourceB)
 {
     FarSkirtVertex polygon[5];uint polygonCount=0;
     FarSkirtClipTriangle(a,b,c,polygon,polygonCount);
@@ -822,6 +1101,14 @@ uint FarSkirtSourceTriangleOutputCount(FarSkirtVertex a,FarSkirtVertex b,FarSkir
     [loop]
     for(uint fan=1;fan+1<polygonCount;++fan)
     {
+        if(StepVoxels<=2)
+        {
+            if(nearRootIndex<FarSkirtNearRootLimit)
+                vertexCount+=FarSkirtNearResolutionTriangleOutputCount(polygon[0],
+                    polygon[fan],polygon[fan+1],nearProbeBase,nearRootIndex,sourceA,sourceB);
+            ++nearRootIndex;
+            continue;
+        }
         const float3 centerVoxel=(polygon[0].position+polygon[fan].position
             +polygon[fan+1].position)/(3.0f*VoxelSizeCm);
         const float centerDensity=FarSkirtDensityAtVoxel(centerVoxel);
@@ -878,9 +1165,15 @@ FarSkirtVertex FarSkirtExtrudedVertexCached(FarSkirtVertex source,float density)
 
 uint FarSkirtEdgeOutputCount(uint edgeVertexId,out FarSkirtCachedSamples cached)
 {
-    cached.endpointDensities=0.0f;
+    if(StepVoxels<=2)
-    cached.centerDensities=0.0f;
+        cached=FarSkirtCachedProbeInput[edgeVertexId];
-    cached.metadata=0;
+    else
+    {
+        cached.endpointDensities=0.0f;
+        cached.centerDensities=0.0f;
+        cached.metadata=0;
+    }
+    if(StepVoxels<=2&&cached.metadata.w!=0u)return 0;
     const uint surfaceVertexCount=min(FarConnectorDrawArgs.Load(OutputArgsByteOffset),OutputVertexCapacity);
     if(edgeVertexId>=surfaceVertexCount)return 0;
     const uint triangleBase=edgeVertexId-(edgeVertexId%3);
@@ -888,11 +1181,21 @@ uint FarSkirtEdgeOutputCount(uint edgeVertexId,out FarSkirtCachedSamples cached)
     const FarSkirtVertex a=FarSkirtSurfaceVertex(triangleBase+edge);
     const FarSkirtVertex b=FarSkirtSurfaceVertex(triangleBase+(edge+1)%3);
     if((FarSkirtBoundaryFaceMask(a.position,b.position)&FarSkirtFaceMask)==0)return 0;
-    const FarSkirtVertex a2=FarSkirtExtrudedVertex(a);
+    FarSkirtVertex a2,b2;
-    const FarSkirtVertex b2=FarSkirtExtrudedVertex(b);
+    if(StepVoxels<=2)
-    cached.endpointDensities=float2(a2.density,b2.density);
+    {
-    return FarSkirtSourceTriangleOutputCount(a,b,b2,cached)
+        a2=FarSkirtExtrudedVertexCached(a,cached.endpointDensities.x);
-        +FarSkirtSourceTriangleOutputCount(a,b2,a2,cached);
+        b2=FarSkirtExtrudedVertexCached(b,cached.endpointDensities.y);
+    }
+    else
+    {
+        a2=FarSkirtExtrudedVertex(a);
+        b2=FarSkirtExtrudedVertex(b);
+        cached.endpointDensities=float2(a2.density,b2.density);
+    }
+    uint nearRootIndex=0;
+    return FarSkirtSourceTriangleOutputCount(a,b,b2,cached,cached.metadata.y,nearRootIndex,a,b)
+        +FarSkirtSourceTriangleOutputCount(a,b2,a2,cached,cached.metadata.y,nearRootIndex,a,b);
 }
 
 void FarSkirtStoreVertex(uint outputId,FarSkirtVertex vertex,uint panelBaseVertex)
@@ -920,9 +1223,40 @@ void FarSkirtEmitTriangle(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex c,
     FarSkirtStoreVertex(edgeBase+edgeCursor++,b,panelBaseVertex);
 }
 
+void FarSkirtEmitNearResolutionTriangle(FarSkirtVertex a,FarSkirtVertex b,
+    FarSkirtVertex c,FarSkirtVertex sourceA,FarSkirtVertex sourceB,uint probeBase,
+    uint rootIndex,uint panelBaseVertex,uint edgeBase,inout uint edgeCursor)
+{
+    if(FarSkirtNearRootMaximum(probeBase,rootIndex)<=0.25f)
+    {
+        FarSkirtEmitTriangle(a,b,c,panelBaseVertex,edgeBase,edgeCursor);
+        return;
+    }
+    uint childFanIndex=0;
+    [unroll]
+    for(uint child=0;child<4;++child)
+    {
+        FarSkirtVertex childA,childB,childC;
+        FarSkirtNearBuildChildTriangle(a,b,c,probeBase,rootIndex,sourceA,sourceB,
+            child,childA,childB,childC);
+        FarSkirtVertex polygon[5];uint polygonCount=0;
+        FarSkirtClipTriangle(childA,childB,childC,polygon,polygonCount);
+        if(polygonCount<3)continue;
+        [loop]
+        for(uint fan=1;fan+1<polygonCount;++fan)
+        {
+            if(FarSkirtNearChildFanOutputCount(probeBase,rootIndex,childFanIndex)!=0)
+                FarSkirtEmitTriangle(polygon[0],polygon[fan],polygon[fan+1],
+                    panelBaseVertex,edgeBase,edgeCursor);
+            ++childFanIndex;
+        }
+    }
+}
+
 void FarSkirtEmitSourceTriangle(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex c,
-    float4 centerDensities,inout uint centerIndex,
+    float4 centerDensities,inout uint centerIndex,uint nearProbeBase,inout uint nearRootIndex,
-    uint panelBaseVertex,uint edgeBase,inout uint edgeCursor)
+    FarSkirtVertex sourceA,FarSkirtVertex sourceB,uint panelBaseVertex,uint edgeBase,
+    inout uint edgeCursor)
 {
     FarSkirtVertex polygon[5];uint polygonCount=0;
     FarSkirtClipTriangle(a,b,c,polygon,polygonCount);
@@ -930,6 +1264,14 @@ void FarSkirtEmitSourceTriangle(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex
     [loop]
     for(uint fan=1;fan+1<polygonCount;++fan)
     {
+        if(StepVoxels<=2)
+        {
+            if(nearRootIndex<FarSkirtNearRootLimit)
+                FarSkirtEmitNearResolutionTriangle(polygon[0],polygon[fan],polygon[fan+1],
+                    sourceA,sourceB,nearProbeBase,nearRootIndex,panelBaseVertex,edgeBase,edgeCursor);
+            ++nearRootIndex;
+            continue;
+        }
         const uint index=centerIndex++;
         const float centerDensity=index==0?centerDensities.x:index==1?centerDensities.y:
             index==2?centerDensities.z:centerDensities.w;
@@ -939,6 +1281,142 @@ void FarSkirtEmitSourceTriangle(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex
     }
 }
 
+[numthreads(64,1,1)]
+void EvaluateFarSkirtEndpointsCS(uint3 dispatchId : SV_DispatchThreadID)
+{
+    const uint edgeVertexId=dispatchId.x;
+    if(edgeVertexId>=OutputVertexCapacity||StepVoxels>2)return;
+    const uint surfaceVertexCount=min(FarConnectorDrawArgs.Load(OutputArgsByteOffset),
+        OutputVertexCapacity);
+    FarSkirtCachedSamples cached;
+    cached.endpointDensities=0.0f;cached.centerDensities=0.0f;cached.metadata=0;
+    if(edgeVertexId<surfaceVertexCount)
+    {
+        const uint triangleBase=edgeVertexId-(edgeVertexId%3);
+        const uint edge=edgeVertexId%3;
+        const FarSkirtVertex a=FarSkirtSurfaceVertex(triangleBase+edge);
+        const FarSkirtVertex b=FarSkirtSurfaceVertex(triangleBase+(edge+1)%3);
+        if((FarSkirtBoundaryFaceMask(a.position,b.position)&FarSkirtFaceMask)!=0)
+        {
+            const FarSkirtVertex a2=FarSkirtExtrudedVertexCached(a,0.0f);
+            const FarSkirtVertex b2=FarSkirtExtrudedVertexCached(b,0.0f);
+            float endpointDensities[2];
+            [loop]
+            for(uint endpoint=0;endpoint<2;++endpoint)
+            {
+                const float3 position=endpoint==0?a2.position:b2.position;
+                endpointDensities[endpoint]=NearSkirtDensityAtVoxel(position/VoxelSizeCm);
+            }
+            cached.endpointDensities=float2(endpointDensities[0],endpointDensities[1]);
+        }
+    }
+    FarSkirtCachedProbeOutput[edgeVertexId]=cached;
+}
+
+[numthreads(64,1,1)]
+void PrepareFarSkirtNearProbesCS(uint3 dispatchId : SV_DispatchThreadID)
+{
+    const uint edgeVertexId=dispatchId.x;
+    if(edgeVertexId>=OutputVertexCapacity||StepVoxels>2)return;
+    FarSkirtCachedSamples cached;
+    cached.endpointDensities=0.0f;cached.centerDensities=0.0f;cached.metadata=0;
+    const uint surfaceVertexCount=min(FarConnectorDrawArgs.Load(OutputArgsByteOffset),
+        OutputVertexCapacity);
+    if(edgeVertexId>=surfaceVertexCount)
+    {
+        FarSkirtCachedProbeOutput[edgeVertexId]=cached;
+        return;
+    }
+    if(edgeVertexId>=FarSkirtNearProbeSurfaceCapacity())
+    {
+        cached.metadata.w=1u;
+        uint previousStatus=0u;
+        FarConnectorDrawArgs.InterlockedOr(OutputArgsByteOffset+40,1u,previousStatus);
+        FarSkirtCachedProbeOutput[edgeVertexId]=cached;
+        return;
+    }
+    const uint triangleBase=edgeVertexId-(edgeVertexId%3);
+    const uint edge=edgeVertexId%3;
+    const FarSkirtVertex a=FarSkirtSurfaceVertex(triangleBase+edge);
+    const FarSkirtVertex b=FarSkirtSurfaceVertex(triangleBase+(edge+1)%3);
+    if((FarSkirtBoundaryFaceMask(a.position,b.position)&FarSkirtFaceMask)==0)
+    {
+        FarSkirtCachedProbeOutput[edgeVertexId]=cached;
+        return;
+    }
+    const FarSkirtCachedSamples endpoints=FarSkirtCachedProbeInput[edgeVertexId];
+    cached.endpointDensities=endpoints.endpointDensities;
+    const FarSkirtVertex a2=FarSkirtExtrudedVertexCached(a,cached.endpointDensities.x);
+    const FarSkirtVertex b2=FarSkirtExtrudedVertexCached(b,cached.endpointDensities.y);
+    const uint rootCount=FarSkirtCountNearSourceTriangleRoots(a,b,b2)
+        +FarSkirtCountNearSourceTriangleRoots(a,b2,a2);
+    uint probeBase=0;
+    if(rootCount!=0)
+        FarSkirtNearProbeCounter.InterlockedAdd(0,rootCount*FarSkirtNearProbeStride,probeBase);
+    uint rootIndex=0;
+    FarSkirtPrepareNearSourceTriangle(a,b,b2,probeBase,rootIndex);
+    FarSkirtPrepareNearSourceTriangle(a,b2,a2,probeBase,rootIndex);
+    cached.metadata.y=probeBase;
+    cached.metadata.z=rootIndex;
+    FarSkirtCachedProbeOutput[edgeVertexId]=cached;
+}
+
+[numthreads(1,1,1)]
+void FinalizeFarSkirtNearProbeDispatchCS(uint3 dispatchId : SV_DispatchThreadID)
+{
+    const uint probeCount=FarSkirtNearProbeCounterInput.Load(0);
+    FarSkirtNearProbeDispatchArgs.Store(0,(probeCount+63u)/64u);
+    FarSkirtNearProbeDispatchArgs.Store(4,1u);
+    FarSkirtNearProbeDispatchArgs.Store(8,1u);
+}
+
+[numthreads(64,1,1)]
+void EvaluateFarSkirtNearProbesCS(uint3 dispatchId : SV_DispatchThreadID)
+{
+    const uint probeIndex=dispatchId.x;
+    if(probeIndex>=FarSkirtNearProbeCounterInput.Load(0))return;
+    if(probeIndex%FarSkirtNearProbeStride>=FarSkirtNearRootProbeStride)return;
+    float4 probe=FarSkirtNearProbeOutput[probeIndex];
+    probe.w=NearSkirtDensityAtVoxel(probe.xyz);
+    FarSkirtNearProbeOutput[probeIndex]=probe;
+}
+
+[numthreads(64,1,1)]
+void PrepareFarSkirtNearChildProbesCS(uint3 dispatchId : SV_DispatchThreadID)
+{
+    const uint edgeVertexId=dispatchId.x;
+    if(edgeVertexId>=OutputVertexCapacity||StepVoxels>2)return;
+    const uint surfaceVertexCount=min(FarConnectorDrawArgs.Load(OutputArgsByteOffset),
+        OutputVertexCapacity);
+    if(edgeVertexId==0u)FarSkirtNearProbeOutput[0].z=(float)surfaceVertexCount;
+    if(edgeVertexId>=surfaceVertexCount)return;
+    const uint triangleBase=edgeVertexId-(edgeVertexId%3);
+    const uint edge=edgeVertexId%3;
+    const FarSkirtVertex a=FarSkirtSurfaceVertex(triangleBase+edge);
+    const FarSkirtVertex b=FarSkirtSurfaceVertex(triangleBase+(edge+1)%3);
+    if((FarSkirtBoundaryFaceMask(a.position,b.position)&FarSkirtFaceMask)==0)return;
+    const FarSkirtCachedSamples cached=FarSkirtCachedProbeInput[edgeVertexId];
+    if(cached.metadata.w!=0u)return;
+    const FarSkirtVertex a2=FarSkirtExtrudedVertexCached(a,cached.endpointDensities.x);
+    const FarSkirtVertex b2=FarSkirtExtrudedVertexCached(b,cached.endpointDensities.y);
+    uint rootIndex=0;
+    FarSkirtPrepareNearChildSourceTriangle(a,b,b2,a,b,cached.metadata.y,rootIndex);
+    FarSkirtPrepareNearChildSourceTriangle(a,b2,a2,a,b,cached.metadata.y,rootIndex);
+}
+
+[numthreads(64,1,1)]
+void EvaluateFarSkirtNearChildProbesCS(uint3 dispatchId : SV_DispatchThreadID)
+{
+    const uint probeIndex=dispatchId.x;
+    if(probeIndex>=FarSkirtNearProbeCounterInput.Load(0))return;
+    const uint probeInRoot=probeIndex%FarSkirtNearProbeStride;
+    if(probeInRoot<FarSkirtNearRootProbeStride)return;
+    float4 probe=FarSkirtNearProbeOutput[probeIndex];
+    if(probe.w< -1.0e30f)return;
+    probe.w=NearSkirtDensityAtVoxel(probe.xyz);
+    FarSkirtNearProbeOutput[probeIndex]=probe;
+}
+
 float FarConnectorRadius(float3 center, float3 direction, float nominalRadius)
 {
     const float3 nominalPosition=center+direction*nominalRadius;
@@ -1097,9 +1575,7 @@ void BuildFarConnectorRingsCS(uint3 groupId : SV_GroupID,
         }
         if (ringId == 0)
         {
-            const uint slot = OutputVertexCapacity > 0
+            const uint panelBaseVertex = FarConnectorOutputBaseVertex;
-                ? OutputBaseVertex / OutputVertexCapacity : 0;
-            const uint panelBaseVertex = slot * FarConnectorPanelVertexCapacity;
             FarConnectorDrawArgs.Store(OutputArgsByteOffset + 20, FarConnectorVertexCount);
             FarConnectorDrawArgs.Store(OutputArgsByteOffset + 24, FarConnectorVertexCount > 0 ? 1 : 0);
             FarConnectorDrawArgs.Store(OutputArgsByteOffset + 28, panelBaseVertex);
@@ -1175,9 +1651,7 @@ void EmitFarConnectorPanelsCS(uint3 dispatchId : SV_DispatchThreadID)
     vertex.normal = ring.normals[cornerIndex].xyz;
     const float connectorLength = length(connector.end - connector.start);
     vertex.uv = float2(ring.t * connectorLength / 256.0f, (float)side / 4.0f);
-    const uint slot = OutputVertexCapacity > 0
+    const uint base = FarConnectorOutputBaseVertex;
-        ? OutputBaseVertex / OutputVertexCapacity : 0;
-    const uint base = slot * FarConnectorPanelVertexCapacity;
     const uint byteAddress = (base + vertexId) * 32;
     FarConnectorVertexBytes.Store3(byteAddress, asuint(vertex.position));
     FarConnectorVertexBytes.Store3(byteAddress + 12, asuint(vertex.normal));
@@ -1382,6 +1856,7 @@ void ClassifyFarSkirtsCS(uint3 dispatchId : SV_DispatchThreadID)
 [numthreads(1,1,1)]
 void ScanFarSkirtsCS(uint3 dispatchId : SV_DispatchThreadID)
 {
+    if(FarConnectorDrawArgs.Load(OutputArgsByteOffset+40)!=0u)return;
     const uint surfaceVertexCount=min(FarConnectorDrawArgs.Load(OutputArgsByteOffset),
         OutputVertexCapacity);
     uint totalSkirtVertices=0;
@@ -1419,14 +1894,18 @@ void EmitFarSkirtsCS(uint3 dispatchId : SV_DispatchThreadID)
     const FarSkirtVertex b=FarSkirtSurfaceVertex(triangleBase+(edge+1)%3);
     if((FarSkirtBoundaryFaceMask(a.position,b.position)&FarSkirtFaceMask)==0)return;
     const FarSkirtCachedSamples cached=FarSkirtCachedProbeInput[edgeVertexId];
+    if(cached.metadata.w!=0u)return;
     const FarSkirtVertex a2=FarSkirtExtrudedVertexCached(a,cached.endpointDensities.x);
     const FarSkirtVertex b2=FarSkirtExtrudedVertexCached(b,cached.endpointDensities.y);
     const uint panelBaseVertex=FarConnectorDrawArgs.Load(OutputArgsByteOffset+28);
     const uint edgeBase=FarConnectorVertexCount+FarSkirtScratch[edgeVertexId];
     uint edgeCursor=0,centerIndex=0;
+    uint nearRootIndex=0;
     FarSkirtEmitSourceTriangle(a,b,b2,cached.centerDensities,centerIndex,
+        cached.metadata.y,nearRootIndex,a,b,
         panelBaseVertex,edgeBase,edgeCursor);
     FarSkirtEmitSourceTriangle(a,b2,a2,cached.centerDensities,centerIndex,
+        cached.metadata.y,nearRootIndex,a,b,
         panelBaseVertex,edgeBase,edgeCursor);
 }
 

Standalone/Source/StandaloneViewer.cpp

@@ -1403,6 +1403,16 @@ namespace
         std::chrono::steady_clock::time_point ConsumedAt{};
         bool bGpuComplete=false,bTimingResolved=false;
     };
+    struct FGpuFarPanelOutputRange
+    {
+        uint32 BaseVertex=0;
+        uint32 CapacityVertices=0;
+    };
+    struct FGpuFarPanelFreeRange
+    {
+        uint32 BaseVertex=0;
+        uint32 CapacityVertices=0;
+    };
 
     enum class EGpuWarmupLaunchPolicy
     {
@@ -1632,7 +1642,9 @@ namespace
         uint32 GpuFarMaxTilesPerSubmissionBatch=0;
         uint32 GpuFarOutputCapacitySlots=0,GpuFarNextOutputSlot=0;
         uint32 GpuFarOutputPoolReservedSlots=0;
-        uint64 GpuFarOutputCountChecks=0,GpuFarOutputOverflowCount=0;
+        uint32 GpuFarPanelOutputHighWater=0,GpuFarPanelOutputReservedCapacityVertices=0;
+        uint64 GpuFarOutputCountChecks=0,GpuFarOutputOverflowCount=0,
+            GpuFarConnectorOverflowCount=0;
         uint64 GpuFarOutputCountReadbackFailures=0;
         uint32 GpuFarMaxRequiredVertices=0,GpuFarMaxRequiredOutputSlot=UINT32_MAX;
         FVoxelTileKey GpuFarMaxRequiredVertexKey{};
@@ -1644,6 +1656,8 @@ namespace
         std::unordered_map<FVoxelTileKey,std::chrono::steady_clock::time_point,FStandaloneTileKeyHash> GpuFarPreparedAt;
         std::vector<FVoxelTileKey> GpuFarCandidateTiles,GpuFarBaseCpuFallbackTiles;
         std::vector<FVoxelTileKey> GpuFarCpuFallbackTiles,GpuFarNeedPreparation;
+        std::unordered_map<FVoxelTileKey,FGpuFarPanelOutputRange,FStandaloneTileKeyHash> GpuFarPanelOutputRanges;
+        std::vector<FGpuFarPanelFreeRange> GpuFarPanelOutputFreeRanges;
         uint64 GpuFarWarmupCpuClaimedTiles=0,GpuFarWarmupGpuTakeoverTiles=0;
         double GpuFarPreparationWallMs=0.0,GpuFarPreparationWorkerMs=0.0;
         uint64 GpuFarPumpCalls=0,GpuFarCompletionPolls=0;
@@ -1664,7 +1678,6 @@ namespace
         bool bGpuFirstBackedPresentReported=false;
         bool bProfileAttribution=false;
         bool bProfileDetailBand=false;
-        bool bGpuFarSkipCpuDetailBand=false;
         bool bCompareEvaluators=false;
         bool bParityOnly=false;
         bool bDensityHash=false;
@@ -1790,6 +1803,8 @@ namespace
             {std::lock_guard Lock(GpuFarPreparedResultsMutex);GpuFarPreparedResults.clear();}
             GpuFarCandidateTiles.clear();GpuFarBaseCpuFallbackTiles.clear();
             GpuFarCpuFallbackTiles.clear();GpuFarNeedPreparation.clear();
+            GpuFarPanelOutputRanges.clear();GpuFarPanelOutputFreeRanges.clear();
+            GpuFarPanelOutputHighWater=GpuFarPanelOutputReservedCapacityVertices=0;
             bGpuFarPreparationStarted=false;bGpuFarInputsFinalized=false;
             bGpuFarPreparationDeferredReported=false;
             bGpuFarPreparationFinished.store(false,std::memory_order_relaxed);
@@ -1811,6 +1826,9 @@ namespace
             GpuFarNearPackedSkippedSamples=GpuFarProtectedRouteSamples=0;
             GpuFarProtectedThinSamples=GpuFarProtectedSolidSamples=0;
             GpuFarProtectedSealSamples=GpuFarProtectedOverlaySamples=0;
+            GpuFarOutputCountChecks=GpuFarOutputOverflowCount=GpuFarConnectorOverflowCount=0;
+            GpuFarOutputCountReadbackFailures=0;GpuFarMaxRequiredVertices=0;
+            GpuFarMaxRequiredOutputSlot=UINT32_MAX;GpuFarMaxRequiredVertexKey={};
             GpuFarGpuDispatchMaxMs=0.0;GpuFarSubmissionSpanMs=0.0;GpuFarSubmissionSpanMaxMs=0.0;
             GpuFarGpuDispatchMaxKey={};bGpuFarGpuDispatchMaxKeyValid=false;
             GpuFarLocalMemoryPeakBytes=0;GpuFarLocalMemoryBudgetBytes=0;
@@ -1821,6 +1839,7 @@ namespace
             GpuFarPoolReserveMs=0.0;
             GpuFarSubmissionBatches=0;GpuFarMaxTilesPerSubmissionBatch=0;
             GpuFarOutputPoolReservedSlots=0;
+            GpuFarPanelOutputReservedCapacityVertices=0;
             DirtyBatchSet.clear();DesiredTiles.clear();DesiredTileSet.clear();CrossingWork.clear();
             DeferredLeaverTiles.clear();
             for(int32 I=0;I<=MaxLevels;++I)
@@ -2103,10 +2122,11 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                         ColdTimelineMilliseconds(PoolStart),PreReservePoolSlots);
                     std::string PoolError;
                     if(VoxelForgeGpuFarViewer::ReserveFarOutputPool(Device,Context,GpuFarShaderPath,
-                        static_cast<uint32>(PreReservePoolSlots),static_cast<uint32>(CoarseTileCells),PoolError))
+                        static_cast<uint32>(PreReservePoolSlots),static_cast<uint32>(CoarseTileCells),PoolError,0))
                     {
                         GpuFarOutputCapacitySlots=static_cast<uint32>(PreReservePoolSlots);
                         GpuFarOutputPoolReservedSlots=GpuFarOutputCapacitySlots;
+                        GpuFarPanelOutputReservedCapacityVertices=0;
                         const double PoolMs=std::chrono::duration<double,std::milli>(
                             std::chrono::steady_clock::now()-PoolStart).count();
                         GpuFarPoolReserveMs+=PoolMs;
@@ -2490,7 +2510,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         void PopulateGpuVisualDetailBandPlan(VoxelForgeCore::ShapeBuildRequest& Request,
             const std::vector<FVoxelTileKey>& TileKeys) const
         {
-            if(!bGpuFarEnabled||(!bProfileDetailBand&&!bGpuFarSkipCpuDetailBand))return;
+            if(!bGpuFarEnabled||!bProfileDetailBand)return;
             std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> TileSet;
             TileSet.reserve(TileKeys.size()*2+1);
             for(const FVoxelTileKey& Key:TileKeys)TileSet.insert(Key);
@@ -2504,11 +2524,10 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 Region.step=std::max(1,Key.ExtentVoxels()/CHUNK_SIZE);
                 const uint8 VisibleSkirts=World->bGenerateSkirts
                     ?ComputeVisibleSkirtFaceMask(Key,TileSet):0;
-                if(VisibleSkirts==0)Request.gpuVisualDetailTiles.push_back(Region);
+                if(VisibleSkirts==0||bGpuSkirtEmitEnabled)
+                    Request.gpuVisualDetailTiles.push_back(Region);
                 else Request.cpuVisualDetailTiles.push_back(Region);
             }
-            Request.skipCpuDetailBandForGpuVisuals=bGpuFarSkipCpuDetailBand
-                &&!Request.gpuVisualDetailTiles.empty();
         }
 
         bool BuildShapeWorldForCamera(std::string& Error,bool bIncludeFarWorld=false,
@@ -2590,7 +2609,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     if(Key.Level<0)continue;
                     if(Key.Level<World->FullResClipLevels)
                     {
-                        if(!World->bGenerateSkirts||ComputeVisibleSkirtFaceMask(Key,PoolDesiredSet)==0)
+                        if(!World->bGenerateSkirts||bGpuSkirtEmitEnabled
+                            ||ComputeVisibleSkirtFaceMask(Key,PoolDesiredSet)==0)
                             ++GpuFarPreReservePoolSlots;
                         continue;
                     }
@@ -3074,10 +3094,12 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                             BatchOutputKeys.push_back(BatchKey);
                         }
                         std::vector<uint32> RequiredVertexCounts;
+                        std::vector<uint32> ConnectorOverflowFlags;
                         std::string CountReadbackError;
                         const bool bOutputCountsRead=
-                            VoxelForgeGpuFarViewer::ReadFarTileRequiredVertexCounts(Device,Context,
+                            VoxelForgeGpuFarViewer::ReadFarTileOutputStatuses(Device,Context,
-                                GpuFarShaderPath,OutputSlots,RequiredVertexCounts,CountReadbackError);
+                                GpuFarShaderPath,OutputSlots,RequiredVertexCounts,
+                                ConnectorOverflowFlags,CountReadbackError);
                         std::vector<FVoxelTileKey> OverflowFallbackKeys;
                         if(!bOutputCountsRead)
                         {
@@ -3112,6 +3134,16 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                                         CheckedTile.OutputSlot,Required,
                                         VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile,CurrentBatchIndex);
                                 }
+                                if(OutputIndex<ConnectorOverflowFlags.size()
+                                    &&ConnectorOverflowFlags[OutputIndex]!=0)
+                                {
+                                    ++GpuFarConnectorOverflowCount;
+                                    if(std::find(OverflowFallbackKeys.begin(),OverflowFallbackKeys.end(),BatchKey)
+                                        ==OverflowFallbackKeys.end())OverflowFallbackKeys.push_back(BatchKey);
+                                    std::fprintf(stderr,"[GpuFarPanelOverflow] key=(%d,%d,%d,lod%d) output_slot=%u batch=%u action=cpu_fallback\n",
+                                        BatchKey.Coord.X,BatchKey.Coord.Y,BatchKey.Coord.Z,BatchKey.Level,
+                                        CheckedTile.OutputSlot,CurrentBatchIndex);
+                                }
                             }
                         }
                         for(const FVoxelTileKey& FallbackKey:OverflowFallbackKeys)
@@ -3120,6 +3152,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                             GpuFarPreparedInputs.erase(FallbackKey);
                             GpuFarPreparationErrors.erase(FallbackKey);
                             GpuFarPreparedAt.erase(FallbackKey);
+                            ReleaseGpuFarPanelOutputRange(FallbackKey);
                             auto FallbackTileIt=GpuFarTiles.find(FallbackKey);
                             if(FallbackTileIt==GpuFarTiles.end())continue;
                             const uint32 ReusableSlot=FallbackTileIt->second.OutputSlot;
@@ -3202,12 +3235,14 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     &&Context->GetData(Tile.CompletionEvent,nullptr,0,D3D11_ASYNC_GETDATA_DONOTFLUSH)==S_OK)
                 {
                     std::vector<uint32> RequiredVertexCounts;
+                    std::vector<uint32> ConnectorOverflowFlags;
                     std::string CountReadbackError;
-                    const bool bOutputCountRead=VoxelForgeGpuFarViewer::ReadFarTileRequiredVertexCounts(
+                    const bool bOutputCountRead=VoxelForgeGpuFarViewer::ReadFarTileOutputStatuses(
                         Device,Context,GpuFarShaderPath,{Tile.OutputSlot},RequiredVertexCounts,
-                        CountReadbackError);
+                        ConnectorOverflowFlags,CountReadbackError);
                     const bool bOverflow=!bOutputCountRead||RequiredVertexCounts.empty()
-                        ||RequiredVertexCounts[0]>VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile;
+                        ||RequiredVertexCounts[0]>VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile
+                        ||ConnectorOverflowFlags.empty()||ConnectorOverflowFlags[0]!=0;
                     if(!bOutputCountRead)
                     {
                         ++GpuFarOutputCountReadbackFailures;
@@ -3231,12 +3266,19 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                                 Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,Tile.OutputSlot,Required,
                                 VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile);
                         }
+                        if(!ConnectorOverflowFlags.empty()&&ConnectorOverflowFlags[0]!=0)
+                        {
+                            ++GpuFarConnectorOverflowCount;
+                            std::fprintf(stderr,"[GpuFarPanelOverflow] key=(%d,%d,%d,lod%d) output_slot=%u action=cpu_fallback\n",
+                                Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,Tile.OutputSlot);
+                        }
                     }
                     if(bOverflow)
                     {
                         const uint32 ReusableSlot=Tile.OutputSlot;
                         GpuFarEligibleTiles.erase(Key);GpuFarPreparedInputs.erase(Key);
                         GpuFarPreparationErrors.erase(Key);GpuFarPreparedAt.erase(Key);
+                        ReleaseGpuFarPanelOutputRange(Key);
                         ReleaseGpuFarTile(Tile);GpuFarTiles.erase(Key);
                         GpuFarFreeOutputSlots.push_back(ReusableSlot);
                         AddGpuFarCpuFallbackTiles({Key});
@@ -3305,15 +3347,14 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 Result.Input=Key.Level<World->FullResClipLevels
                     ?ShapeWorld->exportGpuFieldTile(CoreKey,Result.Error)
                     :ShapeWorld->exportFarFieldTile(CoreKey,Result.Error);
-                if(Result.Input&&Key.Level>=World->FullResClipLevels)
+                if(Result.Input&&VisibleSkirtFaceMask!=0)
                 {
                     Result.Input->visibleSkirtFaceMask=VisibleSkirtFaceMask;
-                    // The CPU far mesher uses extent / CoarseTileCells, which can
-                    // differ from Key.StepVoxels() when coarse tiles use a reduced
-                    // cell count. Reuse the exported mesher step so skirt depth
-                    // stays identical to the CPU reference.
                     Result.Input->skirtDepthVoxels=std::max(1.0f,
                         static_cast<float>(World->SkirtCells))*Result.Input->step;
+                }
+                if(Result.Input&&Key.Level>=World->FullResClipLevels)
+                {
                     if(!bGpuConnectorRepairEnabled)
                     {
                         Result.Input->connectorPanels.clear();
@@ -3382,6 +3423,104 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 Tile.ConnectorVertices);
         }
 
+        void CoalesceGpuFarPanelOutputFreeRanges()
+        {
+            std::sort(GpuFarPanelOutputFreeRanges.begin(),GpuFarPanelOutputFreeRanges.end(),
+                [](const FGpuFarPanelFreeRange& A,const FGpuFarPanelFreeRange& B)
+                    {return A.BaseVertex<B.BaseVertex;});
+            std::vector<FGpuFarPanelFreeRange> Merged;
+            Merged.reserve(GpuFarPanelOutputFreeRanges.size());
+            for(const FGpuFarPanelFreeRange& Range:GpuFarPanelOutputFreeRanges)
+            {
+                if(Range.CapacityVertices==0)continue;
+                if(!Merged.empty()&&static_cast<uint64>(Merged.back().BaseVertex)
+                    +Merged.back().CapacityVertices==Range.BaseVertex)
+                    Merged.back().CapacityVertices+=Range.CapacityVertices;
+                else Merged.push_back(Range);
+            }
+            GpuFarPanelOutputFreeRanges=std::move(Merged);
+            while(!GpuFarPanelOutputFreeRanges.empty())
+            {
+                const FGpuFarPanelFreeRange& Last=GpuFarPanelOutputFreeRanges.back();
+                if(static_cast<uint64>(Last.BaseVertex)+Last.CapacityVertices
+                    !=GpuFarPanelOutputHighWater)break;
+                GpuFarPanelOutputHighWater=Last.BaseVertex;
+                GpuFarPanelOutputFreeRanges.pop_back();
+            }
+        }
+
+        void ReleaseGpuFarPanelOutputRange(const FVoxelTileKey& Key)
+        {
+            const auto It=GpuFarPanelOutputRanges.find(Key);
+            if(It==GpuFarPanelOutputRanges.end())return;
+            if(It->second.CapacityVertices!=0)
+                GpuFarPanelOutputFreeRanges.push_back({It->second.BaseVertex,It->second.CapacityVertices});
+            GpuFarPanelOutputRanges.erase(It);
+            CoalesceGpuFarPanelOutputFreeRanges();
+        }
+
+        bool AllocateGpuFarPanelOutputRange(std::uint32_t CapacityVertices,
+            FGpuFarPanelOutputRange& Out)
+        {
+            if(CapacityVertices==0){Out={};return true;}
+            for(auto It=GpuFarPanelOutputFreeRanges.begin();It!=GpuFarPanelOutputFreeRanges.end();++It)
+            {
+                if(It->CapacityVertices<CapacityVertices)continue;
+                Out={It->BaseVertex,CapacityVertices};
+                It->BaseVertex+=CapacityVertices;
+                It->CapacityVertices-=CapacityVertices;
+                if(It->CapacityVertices==0)GpuFarPanelOutputFreeRanges.erase(It);
+                return true;
+            }
+            if(static_cast<uint64>(GpuFarPanelOutputHighWater)+CapacityVertices
+                >std::numeric_limits<uint32>::max())return false;
+            Out={GpuFarPanelOutputHighWater,CapacityVertices};
+            GpuFarPanelOutputHighWater+=CapacityVertices;
+            return true;
+        }
+
+        bool ReserveGpuFarPanelOutputPool(std::string& Error)
+        {
+            for(auto It=GpuFarPanelOutputRanges.begin();It!=GpuFarPanelOutputRanges.end();)
+            {
+                if(GpuFarEligibleTiles.contains(It->first)){++It;continue;}
+                const FVoxelTileKey Key=It->first;++It;
+                ReleaseGpuFarPanelOutputRange(Key);
+            }
+            for(const FVoxelTileKey& Key:GpuFarQueue)
+            {
+                const auto InputIt=GpuFarPreparedInputs.find(Key);
+                if(InputIt==GpuFarPreparedInputs.end())
+                {Error="prepared input is missing while reserving its panel output";return false;}
+                const VoxelForgeCore::FarFieldTileInput& Input=InputIt->second;
+                const uint64 RequiredCapacity64=static_cast<uint64>(Input.connectorVertexCount)
+                    +(Input.visibleSkirtFaceMask!=0
+                        ?VoxelForgeGpuFarViewer::NearSkirtOutputVertexCapacityPerTile:0u);
+                if(RequiredCapacity64>std::numeric_limits<uint32>::max())
+                {Error="per-tile panel output capacity exceeds the 32-bit vertex limit";return false;}
+                const uint32 RequiredCapacity=static_cast<uint32>(RequiredCapacity64);
+                auto RangeIt=GpuFarPanelOutputRanges.find(Key);
+                if(RangeIt!=GpuFarPanelOutputRanges.end()
+                    &&RangeIt->second.CapacityVertices>=RequiredCapacity)continue;
+                if(RangeIt!=GpuFarPanelOutputRanges.end())ReleaseGpuFarPanelOutputRange(Key);
+                FGpuFarPanelOutputRange Range{};
+                if(!AllocateGpuFarPanelOutputRange(RequiredCapacity,Range))
+                {Error="panel output pool exceeds the 32-bit vertex limit";return false;}
+                GpuFarPanelOutputRanges.emplace(Key,Range);
+            }
+            if(!VoxelForgeGpuFarViewer::ReserveFarOutputPool(Device,Context,GpuFarShaderPath,
+                GpuFarOutputCapacitySlots,static_cast<uint32>(World->CoarseTileCells),Error,
+                GpuFarPanelOutputHighWater))return false;
+            GpuFarOutputPoolReservedSlots=std::max(GpuFarOutputPoolReservedSlots,
+                GpuFarOutputCapacitySlots);
+            GpuFarPanelOutputReservedCapacityVertices=std::max(
+                GpuFarPanelOutputReservedCapacityVertices,GpuFarPanelOutputHighWater);
+            std::printf("[GpuPanelOutputPool] queued_tiles=%zu live_ranges=%zu capacity_vertices=%u high_water=%u status=reserved\n",
+                GpuFarQueue.size(),GpuFarPanelOutputRanges.size(),
+                GpuFarPanelOutputReservedCapacityVertices,GpuFarPanelOutputHighWater);
+            return true;
+        }
+
         void ScheduleGpuFarTiles(const std::vector<FVoxelTileKey>& Keys)
         {
             if(GpuFarPreparationState)
@@ -3405,7 +3544,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld())return;
             std::vector<FVoxelTileKey> CpuTiles;CpuTiles.reserve(Keys.size());
             std::vector<FVoxelTileKey> Eligible;Eligible.reserve(Keys.size());
-            uint64 SkirtTransitionFallbackTiles=0;
+            uint64 NearGpuSkirtTiles=0,SkirtTransitionFallbackTiles=0;
             for(const FVoxelTileKey& Key:Keys)
             {
                 const VoxelForgeCore::TileKey CoreKey{{Key.Coord.X,Key.Coord.Y,Key.Coord.Z},Key.Level};
@@ -3426,14 +3565,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                             GpuFarFreeOutputSlots.push_back(ReusedSlot);
                         }
                     }
-                    if(VisibleSkirtFaces!=0&&NearGpuEligible)
+                    if(VisibleSkirtFaces!=0&&NearGpuEligible&&!bGpuSkirtEmitEnabled)
                     {
-                        // Keep fine tiles on the CPU reference path. Far tiles carry the
+                        // Preserve the CPU route when GPU skirt emission is explicitly disabled.
-                        // CPU-decided visible face mask to the deterministic GPU skirt pass.
                         CpuTiles.push_back(Key);
                         ++SkirtTransitionFallbackTiles;
                     }
-                    else Eligible.push_back(Key);
+                    else
+                    {
+                        if(VisibleSkirtFaces!=0&&NearGpuEligible)++NearGpuSkirtTiles;
+                        Eligible.push_back(Key);
+                    }
                 }
                 else CpuTiles.push_back(Key);
                 if(bDrawOnlyTile&&IsDrawOnlyTile(Key))
@@ -3443,6 +3585,12 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             }
             const std::unordered_set<FVoxelTileKey,FStandaloneTileKeyHash> CurrentEligible(
                 Eligible.begin(),Eligible.end());
+            for(auto It=GpuFarPanelOutputRanges.begin();It!=GpuFarPanelOutputRanges.end();)
+            {
+                if(CurrentEligible.contains(It->first)){++It;continue;}
+                const FVoxelTileKey Key=It->first;++It;
+                ReleaseGpuFarPanelOutputRange(Key);
+            }
             for(auto It=GpuFarTiles.begin();It!=GpuFarTiles.end();)
             {
                 if(CurrentEligible.contains(It->first)){++It;continue;}
@@ -3493,7 +3641,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             std::printf("[GpuFieldEligibility] candidate_gpu_tiles=%zu near_gpu_tiles=%zu cpu_fallback_tiles=%zu warmup_cpu_fallback=%s\n",
                 GpuFarCandidateTiles.size(),NearGpuTiles,GpuFarCpuFallbackTiles.size(),
                 bGpuWarmupRoutePending?"yes":"no");
-            std::printf("[GpuSkirtSeamFallback] near_tiles=%llu reason=fine_visible_lod_skirts gpu_far_skirts=active\n",
+            std::printf("[GpuSkirtSeamRouting] near_skirt_gpu_tiles=%llu near_cpu_fallback_tiles=%llu\n",
+                static_cast<unsigned long long>(NearGpuSkirtTiles),
                 static_cast<unsigned long long>(SkirtTransitionFallbackTiles));
         }
 
@@ -3536,11 +3685,12 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     ColdTimelineMilliseconds(ReserveStart),GpuFarOutputCapacitySlots);
                 std::string PoolError;
                 if(!VoxelForgeGpuFarViewer::ReserveFarOutputPool(Device,Context,GpuFarShaderPath,
-                    GpuFarOutputCapacitySlots,static_cast<uint32>(World->CoarseTileCells),PoolError))
+                    GpuFarOutputCapacitySlots,static_cast<uint32>(World->CoarseTileCells),PoolError,0))
                 {
                     std::fprintf(stderr,"[GpuFarPoolError] %s\n",PoolError.c_str());
                     bGpuFarEnabled=false;bGpuFarPreparationStarted=true;
                     GpuFarOutputCapacitySlots=0;GpuFarOutputPoolReservedSlots=0;
+                    GpuFarPanelOutputReservedCapacityVertices=0;
                     for(auto& Pair:GpuFarTiles)ReleaseGpuFarTile(Pair.second);
                     GpuFarTiles.clear();GpuFarCandidateTiles.clear();GpuFarQueue.clear();GpuFarQueued.clear();
                     GpuFarNeedPreparation.clear();GpuFarEligibleTiles.clear();
@@ -3606,8 +3756,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             std::vector<std::function<void()>> Tasks;Tasks.reserve(NeedPreparation.size());
             for(const FVoxelTileKey& Key:NeedPreparation)
             {
-                const uint8 VisibleSkirtFaceMask=Key.Level>=World->FullResClipLevels
+                const uint8 VisibleSkirtFaceMask=World->bGenerateSkirts&&bGpuSkirtEmitEnabled
-                    &&World->bGenerateSkirts&&bGpuSkirtEmitEnabled
                     ?ComputeVisibleSkirtFaceMask(Key,DesiredTileSet):0;
                 Tasks.emplace_back([this,Key,VisibleSkirtFaceMask,Batch]()
                     {PrepareGpuFarInputTile(Key,VisibleSkirtFaceMask,Batch);});
@@ -3819,6 +3968,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             }
             if(!GpuFarQueue.empty())SortGpuFarQueue();
             AddGpuFarCpuFallbackTiles(Failed);
+            std::string PanelPoolError;
+            if(!ReserveGpuFarPanelOutputPool(PanelPoolError))
+            {
+                std::fprintf(stderr,"[GpuFarPoolError] panel-output reservation failed: %s\n",
+                    PanelPoolError.c_str());
+                bGpuFarEnabled=false;GpuFarQueue.clear();GpuFarQueued.clear();
+                GpuFarEligibleTiles.clear();
+                AddGpuFarCpuFallbackTiles(DesiredTiles);
+                bGpuFarInputsFinalized=true;
+                return false;
+            }
             bGpuFarInputsFinalized=true;
             const std::size_t NearGpuTiles=static_cast<std::size_t>(std::count_if(
                 GpuFarEligibleTiles.begin(),GpuFarEligibleTiles.end(),[&](const FVoxelTileKey& Key)
@@ -3868,7 +4028,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             StartGpuFarPreparation();
             if(!CompleteGpuShaderWarmupAndRoute())return;
             StartGpuFarPreparation();
-            if(!bGpuFarEnabled||!FinalizeGpuFarPreparation()||GpuFarQueue.empty())return;
+            if(!bGpuFarEnabled||!FinalizeGpuFarPreparation()||!bGpuFarInputsFinalized
+                ||GpuFarQueue.empty())return;
             const bool bColdDrainBeforeWholeView=!bInitialClipmapReported&&!bGpuFarAuditEnabled
                 &&!bDisableGpuColdStartDrain&&bGpuFarWarmupBeforeWorkers;
             constexpr std::size_t StartupFirstGpuBatchTarget=32;
@@ -3954,6 +4115,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 Tile.ConnectorRings=PreparedInput->connectorRingCount;
                 Tile.ConnectorVertices=PreparedInput->connectorVertexCount;
                 Tile.VisibleSkirtFaceMask=PreparedSkirtFaceMask;
+                const auto PanelRangeIt=GpuFarPanelOutputRanges.find(Key);
+                if(PanelRangeIt==GpuFarPanelOutputRanges.end())
+                {
+                    GpuFarFreeOutputSlots.push_back(Tile.OutputSlot);
+                    std::fprintf(stderr,"[GpuFarPoolError] panel-output range is missing for key=(%d,%d,%d,lod%d)\n",
+                        Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level);
+                    GpuFarPanelOutputRanges.erase(Key);
+                    bGpuFarEnabled=false;
+                    break;
+                }
+                const FGpuFarPanelOutputRange PanelRange=PanelRangeIt->second;
                 if(!VoxelForgeGpuFarViewer::BuildFarTile(Device,Context,*ShapeWorld,CoreKey,
                     PreparedInput,
                     GpuFarShaderPath,&Tile.VertexBuffer,&Tile.IndexBuffer,&Tile.IndirectArgs,
@@ -3962,9 +4134,11 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     bDeferredBatch?nullptr:&Tile.CompletionEvent,
                     &Tile.TimingQueries,
                     ((bGpuFarAuditEnabled&&Key.Level>=6)||bAuditDrawOnlyGpuTile)?&AuditMesh:nullptr,
-                    &EnqueueMs,&GpuDispatchMs,Error))
+                    &EnqueueMs,&GpuDispatchMs,Error,true,
+                    PanelRange.BaseVertex,PanelRange.CapacityVertices))
                 {
                     GpuFarFreeOutputSlots.push_back(Tile.OutputSlot);
+                    ReleaseGpuFarPanelOutputRange(Key);
                     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());
                     if(PreparedInputIt!=GpuFarPreparedInputs.end())GpuFarPreparedInputs.erase(PreparedInputIt);
@@ -4707,7 +4881,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 std::string PageError;
                 size_t CompiledPageCount=0;
                 if(!ShapeWorld->compileTileRegion(NewCoreTiles,
-                    static_cast<uint32>(std::clamp(WorkerCount,1,8)),CoreScheduler,PageError,
+                    static_cast<uint32>(std::clamp(WorkerCount,1,16)),CoreScheduler,PageError,
                     &CompiledPageCount,&PageStats,true))
                 {
                     std::fprintf(stderr,"[ShapeRegionPage] compile failed at center=(%d,%d,%d): %s\n",
@@ -4729,7 +4903,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     NewDesired.CenterChunk.X,NewDesired.CenterChunk.Y,NewDesired.CenterChunk.Z,
                     NewCoreTiles.size(),static_cast<unsigned long long>(PageCoveredTiles),
                     PageStats.regionsBuilt,static_cast<unsigned long long>(PageStats.latticeNodes),
-                    std::clamp(WorkerCount,1,8),PageStats.shapeSeconds,PageStats.intervalSeconds,
+                    std::clamp(WorkerCount,1,16),PageStats.shapeSeconds,PageStats.intervalSeconds,
                     PageStats.surfaceSeconds,PageStats.workerSeconds,PageSeconds,
                     ResidentTiles.size()-std::min(ResidentTiles.size(),Leavers.size()),Leavers.size());
             }
@@ -5268,9 +5442,10 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 }
                 ReportTileTimings();
                 if(GpuFarOutputCountChecks>0)
-                    std::printf("[GpuFarSlotCapacitySummary] checked_tiles=%llu overflow_tiles=%llu count_readback_failures=%llu capacity_vertices=%u max_required_vertices=%u max_key=(%d,%d,%d,lod%d) max_slot=%u check_path=normal_after_completion_event\n",
+                    std::printf("[GpuFarSlotCapacitySummary] checked_tiles=%llu surface_overflow_tiles=%llu panel_overflow_tiles=%llu count_readback_failures=%llu capacity_vertices=%u max_required_vertices=%u max_key=(%d,%d,%d,lod%d) max_slot=%u check_path=normal_after_completion_event\n",
                         static_cast<unsigned long long>(GpuFarOutputCountChecks),
                         static_cast<unsigned long long>(GpuFarOutputOverflowCount),
+                        static_cast<unsigned long long>(GpuFarConnectorOverflowCount),
                         static_cast<unsigned long long>(GpuFarOutputCountReadbackFailures),
                         VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile,
                         GpuFarMaxRequiredVertices,GpuFarMaxRequiredVertexKey.Coord.X,
@@ -6292,7 +6467,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
 
         std::atomic<size_t> Next{0};
         std::vector<std::thread> Threads;
-        const int32 ThreadCount=std::clamp(Workers,1,8);
+        const int32 ThreadCount=std::clamp(Workers,1,16);
         Threads.reserve(ThreadCount);
         for(int32 Worker=0;Worker<ThreadCount;++Worker)
         {
@@ -6691,7 +6866,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         ShapeRequest.maxVoxel={FineMax.X,FineMax.Y,FineMax.Z};
         ShapeRequest.clipMinVoxel={ClipMin.X,ClipMin.Y,ClipMin.Z};
         ShapeRequest.clipMaxVoxel={ClipMax.X,ClipMax.Y,ClipMax.Z};
-        ShapeRequest.workerLimit=static_cast<uint32>(std::clamp(WorkerCount,1,8));
+        ShapeRequest.workerLimit=static_cast<uint32>(std::clamp(WorkerCount,1,16));
         ShapeRequest.detail=DetailVariant;ShapeRequest.includeFarWorld=bIncludeFarWorld;
         ShapeRequest.roundTunnels=bRoundTunnels;
         ShapeRequest.enhancedTunnelWalls=bEnhancedTunnelWalls;
@@ -7377,7 +7552,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         request.maxVoxel={fineMax.X,fineMax.Y,fineMax.Z};
         request.clipMinVoxel={clipMin.X,clipMin.Y,clipMin.Z};
         request.clipMaxVoxel={clipMax.X,clipMax.Y,clipMax.Z};
-        request.workerLimit=static_cast<uint32>(std::clamp(workers,1,8));
+        request.workerLimit=static_cast<uint32>(std::clamp(workers,1,16));
         request.detail=detail;request.roundTunnels=roundTunnels;
         request.enhancedTunnelWalls=enhancedWalls;request.includeFarWorld=false;
         return true;
@@ -7409,7 +7584,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         std::string error;
         const std::array<VoxelForgeCore::TileKey,1> pageTiles={target};
         std::size_t pageTilesA=0;
-        if(!base.compileTileRegion(pageTiles,static_cast<uint32>(std::clamp(workers,1,8)),
+        if(!base.compileTileRegion(pageTiles,static_cast<uint32>(std::clamp(workers,1,16)),
             scheduler,error,&pageTilesA,nullptr,false))
         {std::fprintf(stderr,"[TilePathHashAudit] path-A page compile failed: %s\n",error.c_str());return 11;}
         auto pathA=base.generateTile(target,options,error);
@@ -7459,9 +7634,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         {std::fprintf(stderr,"[TilePathHashDirections] setup failed: %s\n",error.c_str());return 11;}
         std::size_t analyticForwardCovered=0,analyticReverseCovered=0;
         VoxelForgeCore::TileRegionCompileStats forwardStats,reverseStats;
-        if(!analyticForward.compileTileRegion(pageTiles,static_cast<uint32>(std::clamp(workers,1,8)),
+        if(!analyticForward.compileTileRegion(pageTiles,static_cast<uint32>(std::clamp(workers,1,16)),
                 scheduler,error,&analyticForwardCovered,&forwardStats,true)
-            ||!analyticReverse.compileTileRegion(pageTiles,static_cast<uint32>(std::clamp(workers,1,8)),
+            ||!analyticReverse.compileTileRegion(pageTiles,static_cast<uint32>(std::clamp(workers,1,16)),
                 scheduler,error,&analyticReverseCovered,&reverseStats,true))
         {std::fprintf(stderr,"[TilePathHashDirections] region compile failed: %s\n",error.c_str());return 11;}
         auto analyticPathForward=analyticForward.generateTile(target,options,error);
@@ -7570,7 +7745,6 @@ int main(int Argc,char** Argv)
         else if(Arg=="--gpu-connectors-disabled")App.bGpuConnectorRepairEnabled=false;
         else if(Arg=="--profile-attribution")App.bProfileAttribution=true;
         else if(Arg=="--profile-detail-band")bProfileDetailBand=true;
-        else if(Arg=="--gpu-far-skip-cpu-detail-band")App.bGpuFarSkipCpuDetailBand=true;
         else if(Arg=="--compare-evaluators")App.bCompareEvaluators=true;
         else if(Arg=="--parity-only")App.bParityOnly=true;
         else if(Arg=="--density-hash")App.bDensityHash=true;
@@ -7685,7 +7859,7 @@ int main(int Argc,char** Argv)
         else if(Arg.starts_with("--workers="))
         {
             const int32 Parsed=std::atoi(std::string(Arg.substr(10)).c_str());
-            App.WorkerCount=std::clamp(Parsed,1,8);
+            App.WorkerCount=std::clamp(Parsed,1,16);
         }
         else if(Arg.starts_with("--clip-radius="))
             App.ClipRadius=std::clamp(std::atoi(std::string(Arg.substr(14)).c_str()),1,64);

Standalone/Source/VoxelForge/Private/VoxelMarchingCubesMesher.cpp

@@ -1783,7 +1783,10 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
                         }
                         if(MaximumProbeDensity>0.25f)
                         {
-                            if(RefinementDepth>=1)continue;
+                            if(RefinementDepth>=1)
+                            {
+                                continue;
+                            }
                             const auto MakeMidpoint=[&](const FSkirtVertex& X,
                                 const FSkirtVertex& Y)
                             {
@@ -1842,9 +1845,9 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
             }
             MeshData.NumSkirtTriangles += (Tris.Num() - BaseTriNum) / 3;
         };
-        EmitSkirts(GroundTris,GroundSkirtFaceMasks);
+            EmitSkirts(GroundTris,GroundSkirtFaceMasks);
-        EmitSkirts(CapTris,CapSkirtFaceMasks);
+            EmitSkirts(CapTris,CapSkirtFaceMasks);
-    }
+        }
 
     // F17 — concatène sol PUIS sky-cap : un run d'indices contigu par polygroup (voir
     // FVoxelMeshData::NumCeilingTriangles). Mêmes triangles qu'avant, seul l'ordre change.