wip(standalone): GPU step 6 -- near field on GPU, exact; cold start still slower than CPU

Near LOD0/1 (detail band, relief, root refinement, normals) now evaluated and
meshed on the GPU. 960 near tiles: zero sign/topology differences vs CPU and
WARP, vertex max 0.0049 mm (was 15 cm); far survival audit passes; CPU hashes
44/44 unchanged. Far connector-panel outlier (64.56 cm, LOD8) remains a
tessellation difference, and remote captures show missing far geometry on GPU.

Cold readiness (viewer clock, my interleaved runs): CPU 2.51/2.49/2.40 s,
GPU 5.05/2.68/2.80 s. The ~2.1 s Present stall hit only the first GPU launch.
Viewer spends ~1.13 s before shape compile starts (startup checks), which is
the gap between the 0.95 s benchmark and ~2.4 s on-screen.

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

Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h

@@ -204,10 +204,25 @@ namespace VoxelForgeCore
     };
     static_assert(sizeof(FarFieldSpatialBin) == 24);
 
+    // Layer-local detail parameters for the shared coherent near field. This record stays
+    // float-only so the same values bind to an HLSL StructuredBuffer.
+    struct FarFieldDetailLayerGpu
+    {
+        std::int32_t lo = 0, hi = -1;
+        float surfaceRoughness = 0.0f, roughnessFrequency = 0.0f;
+        float verticalScale = 1.0f, domainWarpStrength = 0.0f;
+        float domainWarpFrequency = 0.0f;
+        std::uint32_t roughnessNoiseType = 0;
+    };
+    static_assert(sizeof(FarFieldDetailLayerGpu) == 32);
+
     struct FarFieldWorldInput
     {
         std::uint32_t spatialBinSize = 32, spatialBinMask = 0;
+        std::uint32_t fieldSeed = 0;
+        std::int32_t detailMinVoxel[3]{}, detailMaxVoxel[3]{};
         std::vector<FarFieldPrimitiveGpu> shapes;
+        std::vector<FarFieldDetailLayerGpu> detailLayers;
         // Open-addressed table. count == 0 denotes an empty slot; mask is capacity - 1.
         std::vector<FarFieldSpatialBin> spatialBins;
         std::vector<std::uint32_t> spatialBinShapeIds;
@@ -236,6 +251,7 @@ namespace VoxelForgeCore
     {
         Vec3i origin{};
         std::int32_t cells = 0, step = 0, sampleDim = 0;
+        bool nearDetailEnabled = false;
         float strataLo = 0.0f, strataHi = 0.0f, strataBlend = 12.0f;
         std::shared_ptr<const FarFieldWorldInput> worldInput;
         // Sparse exception bitset: bit 1 forces route-protected air. Ordinary rock/air
@@ -244,6 +260,11 @@ namespace VoxelForgeCore
         // Sparse exception bitset for exact solid overlays/thin route-adjacent features.
         std::vector<std::uint32_t> protectedSolid;
         std::vector<std::uint32_t> protectedThin;
+        // Near tiles mirror the CPU's sparse interpolation proxy cache. UINT_MAX
+        // means analytic-only; zero means the packed fixed fallback; low values
+        // store |SurfaceDistanceQ|. For Q==1 only, bit 30 marks the ambiguous
+        // packed-sign exception and bit 31 carries its CPU rock sign.
+        std::vector<std::uint32_t> nearPackedProxy;
         std::vector<FarFieldStratumBoundary> stratumBoundaries;
         std::vector<FarFieldConnectorPanel> connectorPanels;
         std::uint32_t connectorRingCount = 0, connectorVertexCount = 0;
@@ -343,6 +364,8 @@ namespace VoxelForgeCore
         std::vector<std::uint32_t> farConnectorCandidates(Vec3i minVoxel, Vec3i maxVoxel) const;
         std::optional<FarFieldTileInput> exportFarFieldTile(const TileKey& key,
                                                              std::string& error) const;
+        std::optional<FarFieldTileInput> exportGpuFieldTile(const TileKey& key,
+                                                             std::string& error) const;
         bool isFarFieldTileEligible(const TileKey& key) const;
 
         bool auditDetailPassing() const;

Standalone/Core/src/ShapeWorldInternal.h

@@ -526,9 +526,20 @@ public:
         auto Data=std::make_shared<VoxelForgeCore::FarFieldWorldInput>();
         Data->spatialBinSize=static_cast<uint32>(ActiveSpatialBinSize);
         Data->spatialBinMask=static_cast<uint32>(TableCapacity-1);
+        Data->fieldSeed=static_cast<uint32>(Owner?Owner->Seed:0);
+        for(int32 Axis=0;Axis<3;++Axis)
+        {Data->detailMinVoxel[Axis]=Min[Axis];Data->detailMaxVoxel[Axis]=Max[Axis];}
         Data->shapes.reserve(SharedFieldPrimitives.size());
         for(const VF_FieldPrimitive& Shape:SharedFieldPrimitives)
             Data->shapes.push_back(ToFarFieldGpuPrimitive(Shape));
+        Data->detailLayers.reserve(LayerFields.size());
+        for(const FLayerField& Field:LayerFields)
+        {
+            const FStrateGenerationParams& Params=Field.Params;
+            Data->detailLayers.push_back({Field.Lo,Field.Hi,Params.SurfaceRoughness,
+                Params.RoughnessFrequency,Params.VerticalScale,Params.DomainWarpStrength,
+                Params.DomainWarpFrequency,static_cast<uint32>(Params.RoughnessNoiseType)});
+        }
         Data->spatialBins.resize(TableCapacity);
         size_t ReferenceCount=0;
         for(const auto& Bin:SpatialBins)ReferenceCount+=Bin.second.size();
@@ -551,20 +562,22 @@ public:
         const double UploadMiB=(G.shapes.size()*sizeof(VoxelForgeCore::FarFieldPrimitiveGpu)
             +G.spatialBins.size()*sizeof(VoxelForgeCore::FarFieldSpatialBin)
             +G.spatialBinShapeIds.size()*sizeof(uint32))/1048576.0;
-        std::printf("[FarGpuWorldData] shapes=%zu spatial_bins=%zu hash_slots=%zu bin_shape_ids=%zu bin_size=%d bytes=%.3fMiB upload_scope=once_per_world\n",
+        std::printf("[FarGpuWorldData] shapes=%zu spatial_bins=%zu hash_slots=%zu bin_shape_ids=%zu bin_size=%d detail_min=(%d,%d,%d) detail_max=(%d,%d,%d) bytes=%.3fMiB upload_scope=once_per_world\n",
             G.shapes.size(),SpatialBins.size(),G.spatialBins.size(),G.spatialBinShapeIds.size(),
-            ActiveSpatialBinSize,UploadMiB);
+            ActiveSpatialBinSize,G.detailMinVoxel[0],G.detailMinVoxel[1],G.detailMinVoxel[2],
+            G.detailMaxVoxel[0],G.detailMaxVoxel[1],G.detailMaxVoxel[2],UploadMiB);
         return CachedFarFieldGpuInput;
     }
 
     bool ExportFarFieldTileInput(const FIntVector& Origin,int32 Cells,int32 Step,
-        VoxelForgeCore::FarFieldTileInput& Out,std::string& Error) const
+        bool bNearDetail,VoxelForgeCore::FarFieldTileInput& Out,std::string& Error) const
     {
         if(!FarBuilt||Cells<=0||Cells>32||Step<=0)
         {Error="far field export requires a built far world and a bounded coarse tile";return false;}
         const int32 Dim=Cells+3;
         const size_t SampleCount=static_cast<size_t>(Dim)*Dim*Dim;
         Out={};Error.clear();Out.origin={Origin.X,Origin.Y,Origin.Z};Out.cells=Cells;Out.step=Step;Out.sampleDim=Dim;
+        Out.nearDetailEnabled=bNearDetail;
         Out.worldInput=GetFarFieldWorldInput(Error);
         if(!Out.worldInput)return false;
         if(LayerRanges.empty())
@@ -596,6 +609,8 @@ public:
         Out.protectedAir.assign((SampleCount+31u)/32u,0u);
         Out.protectedSolid.assign((SampleCount+31u)/32u,0u);
         Out.protectedThin.assign((SampleCount+31u)/32u,0u);
+        if(bNearDetail)
+            Out.nearPackedProxy.assign(SampleCount,std::numeric_limits<uint32>::max());
 
         const int32 LatticeMin[3]={Origin.X-Step,Origin.Y-Step,Origin.Z-Step};
         const int32 LatticeMax[3]={Origin.X+(Cells+1)*Step,Origin.Y+(Cells+1)*Step,
@@ -723,7 +738,34 @@ public:
         for(int32 X=-1;X<=Cells+1;++X,++SampleIndex)
         {
             const int32 WX=Origin.X+X*Step,WY=Origin.Y+Y*Step,WZ=Origin.Z+Z*Step;
-            if(WX>=Min[0]&&WX<=Max[0]&&WY>=Min[1]&&WY<=Max[1]&&WZ>=Min[2]&&WZ<=Max[2])
+            if(bNearDetail&&!bAnalyticOnly&&Contains(WX,WY,WZ))
+            {
+                const size_t Column=static_cast<size_t>(WY-Min[1])*Size[0]
+                    +static_cast<size_t>(WX-Min[0]);
+                const size_t Zi=static_cast<size_t>(WZ-Min[2]);
+                const size_t Word=Column*WordsPerZ+(Zi>>6);
+                const uint64 Mask=uint64{1}<<(Zi&63u);
+                if((CandidateWords[Word]&Mask)==0)
+                    Out.nearPackedProxy[SampleIndex]=0u;
+                else
+                {
+                    const size_t Node=Zi*static_cast<size_t>(Size[1])*Size[0]+Column;
+                    const int32 Q=std::abs(static_cast<int32>(SurfaceDistanceQ[Node]));
+                    const uint32 Magnitude=static_cast<uint32>(std::clamp(Q,1,32767));
+                    if(Magnitude==1u)
+                    {
+                        // At the minimum quantization step a one-ULP GPU sign flip
+                        // can change an entire marching-cubes cell. Keep this sparse
+                        // ambiguous boundary decision aligned with the packed cache.
+                        const bool bRock=(RockWords[Word]&Mask)!=0;
+                        Out.nearPackedProxy[SampleIndex]=0x40000000u
+                            |(bRock?0x80000000u:0u)|Magnitude;
+                    }
+                    else Out.nearPackedProxy[SampleIndex]=Magnitude;
+                }
+            }
+            if(!bNearDetail&&WX>=Min[0]&&WX<=Max[0]&&WY>=Min[1]&&WY<=Max[1]
+                &&WZ>=Min[2]&&WZ<=Max[2])
             {Error="tile lattice intersects the near analytic/detail region; this export is far-field only";return false;}
             if(InsideCritical(WX,WY,WZ))
             {
@@ -1746,6 +1788,7 @@ private:
         if(P.bCurvedFloorClip)Out.flags|=VF_FIELD_PRIMITIVE_CURVED_FLOOR;
         if(P.bCarve)Out.flags|=VF_FIELD_PRIMITIVE_CARVE;
         if(P.bCritical)Out.flags|=VF_FIELD_PRIMITIVE_CRITICAL;
+        if(P.bPassageSurface)Out.flags|=VF_FIELD_PRIMITIVE_PASSAGE;
         for(int32 A=0;A<3;++A)
         {
             Out.minVoxel[A]=P.Min[A];Out.maxVoxel[A]=P.Max[A];
@@ -4081,68 +4124,12 @@ private:
     float OwnerCaveRoughnessCore(double X,double Y,double Z,
                                  const FStrateGenerationParams& Params) const
     {
-        const uint32 Seed=static_cast<uint32>(Owner->Seed);
+        // Keep CPU prediction and the near-field HLSL on one scalar source. The helper
-        const float RF=Params.RoughnessFrequency;
+        // preserves the legacy seed offsets, octave order, and roughness families.
-        const float NoiseX=static_cast<float>(X);
+        return VF_FieldOwnerRoughnessCore(static_cast<float>(X),static_cast<float>(Y),
-        const float NoiseY=static_cast<float>(Y);
+            static_cast<float>(Z),static_cast<uint32>(Owner->Seed),Params.SurfaceRoughness,
-        const float EffectiveZ=(Params.VerticalScale>0.0f&&Params.VerticalScale!=1.0f)
+            Params.RoughnessFrequency,Params.VerticalScale,Params.DomainWarpStrength,
-            ?static_cast<float>(Z)/Params.VerticalScale:static_cast<float>(Z);
+            Params.DomainWarpFrequency,static_cast<uint32>(Params.RoughnessNoiseType));
-        FVector3f Main(NoiseX*RF+VoxelHash::SeedOffset(Seed,11.3f),
-            NoiseY*RF+VoxelHash::SeedOffset(Seed,13.7f),
-            EffectiveZ*RF+VoxelHash::SeedOffset(Seed,17.1f));
-        FVector3f Fine(NoiseX*RF*3.0f+VoxelHash::SeedOffset(Seed,19.1f)+2000.0f,
-            NoiseY*RF*3.0f+VoxelHash::SeedOffset(Seed,23.7f)+2500.0f,
-            EffectiveZ*RF*3.0f+VoxelHash::SeedOffset(Seed,29.3f)+3000.0f);
-        if(Params.DomainWarpStrength>0.0f)
-        {
-            const float WF=Params.DomainWarpFrequency,WS=Params.DomainWarpStrength;
-            float WX=0.0f,WY=0.0f,WZ=0.0f;
-            VoxelNoise::Perlin3D_x3(
-                NoiseX*WF+VoxelHash::SeedOffset(Seed,5.2f),
-                NoiseY*WF+VoxelHash::SeedOffset(Seed,1.3f),
-                EffectiveZ*WF+VoxelHash::SeedOffset(Seed,9.7f),
-                NoiseX*WF+100.0f+VoxelHash::SeedOffset(Seed,7.7f),
-                NoiseY*WF+200.0f+VoxelHash::SeedOffset(Seed,3.1f),
-                EffectiveZ*WF+300.0f,
-                NoiseX*WF+400.0f,
-                NoiseY*WF+500.0f+VoxelHash::SeedOffset(Seed,11.9f),
-                EffectiveZ*WF+600.0f+VoxelHash::SeedOffset(Seed,13.3f),WX,WY,WZ);
-            WX*=VOXEL_NOISE_SCALE*WS;WY*=VOXEL_NOISE_SCALE*WS;WZ*=VOXEL_NOISE_SCALE*WS;
-            Main+=FVector3f(WX,WY,WZ);Fine+=FVector3f(WX,WY,WZ);
-        }
-        float Rough=0.0f,FineValue=0.0f;
-        switch(Params.RoughnessNoiseType)
-        {
-        case EVoxelNoiseType::Ridged:
-        {
-            float MainFBM=0.0f,MainRidged=0.0f,FineFBM=0.0f,FineRidged=0.0f;
-            VoxelNoise::FractalPair<false,true>(Main.X,Main.Y,Main.Z,3,
-                Fine.X,Fine.Y,Fine.Z,2,MainFBM,MainRidged,FineFBM,FineRidged);
-            Rough=MainRidged;FineValue=FineRidged;
-            break;
-        }
-        case EVoxelNoiseType::Mixed:
-        {
-            float MainFBM=0.0f,MainRidged=0.0f,FineFBM=0.0f,FineRidged=0.0f;
-            VoxelNoise::FractalPair<true,true>(Main.X,Main.Y,Main.Z,3,
-                Fine.X,Fine.Y,Fine.Z,2,MainFBM,MainRidged,FineFBM,FineRidged);
-            Rough=MainFBM*0.5f+MainRidged*0.5f;
-            FineValue=FineFBM*0.5f+FineRidged*0.5f;
-            break;
-        }
-        case EVoxelNoiseType::Cellular:
-            Rough=VoxelNoise::Cellular3D(Main);FineValue=VoxelNoise::Cellular3D(Fine);break;
-        case EVoxelNoiseType::FBM:
-        default:
-        {
-            float MainFBM=0.0f,MainRidged=0.0f,FineFBM=0.0f,FineRidged=0.0f;
-            VoxelNoise::FractalPair<true,false>(Main.X,Main.Y,Main.Z,3,
-                Fine.X,Fine.Y,Fine.Z,2,MainFBM,MainRidged,FineFBM,FineRidged);
-            Rough=MainFBM;FineValue=FineFBM;
-            break;
-        }
-        }
-        return (Rough+FineValue*0.4f)*VOXEL_NOISE_SCALE*Params.SurfaceRoughness;
     }
 
     float OwnerCaveRoughness(double X,double Y,double Z,float SurfaceSDF,

Standalone/Core/src/VoxelForgeCore.cpp

@@ -746,11 +746,30 @@ namespace VoxelForgeCore
         const int32 cells=std::max(1,impl_->owner->CoarseTileCells);
         const int32 step=std::max(1,extent/cells);
         FarFieldTileInput result;
-        if(!impl_->shape->ExportFarFieldTileInput(ToNative(key).OriginVoxels(),cells,step,result,error))
+        if(!impl_->shape->ExportFarFieldTileInput(ToNative(key).OriginVoxels(),cells,step,false,result,error))
             return std::nullopt;
         return result;
     }
 
+    std::optional<FarFieldTileInput> World::exportGpuFieldTile(const TileKey& key,
+                                                                std::string& error) const
+    {
+        if(!impl_||!impl_->owner||!impl_->shape||!impl_->shape->IsBuilt())
+        {error="world must be compiled before exporting a GPU field tile";return std::nullopt;}
+        if(key.lod<0||key.lod>8)
+        {error="GPU field export accepts LOD0 through LOD8";return std::nullopt;}
+        if(!impl_->shape->HasFarWorld())
+        {error="world must include the far parent/connector shapes before export";return std::nullopt;}
+        const int32 extent=CHUNK_SIZE<<key.lod;
+        const bool nearDetail=key.lod<impl_->configuration.fullResolutionLodCount;
+        const int32 cells=nearDetail?CHUNK_SIZE:std::max(1,impl_->owner->CoarseTileCells);
+        const int32 step=std::max(1,extent/cells);
+        FarFieldTileInput result;
+        if(!impl_->shape->ExportFarFieldTileInput(ToNative(key).OriginVoxels(),cells,step,
+            nearDetail,result,error))return std::nullopt;
+        return result;
+    }
+
     bool World::isFarFieldTileEligible(const TileKey& key) const
     {
         if(!impl_||!impl_->owner||!impl_->shape||!impl_->shape->IsBuilt()

Standalone/Core/src/VoxelForgeFieldShared.h

@@ -68,6 +68,7 @@ static const VF_FieldI32 VF_FIELD_PRIMITIVE_HAS_FLOOR = 8;
 static const VF_FieldI32 VF_FIELD_PRIMITIVE_CURVED_FLOOR = 16;
 static const VF_FieldI32 VF_FIELD_PRIMITIVE_CARVE = 32;
 static const VF_FieldI32 VF_FIELD_PRIMITIVE_CRITICAL = 64;
+static const VF_FieldI32 VF_FIELD_PRIMITIVE_PASSAGE = 128;
 
 VF_FIELD_INLINE VF_FieldReal VF_FieldEvaluateFloorProfile(VF_FieldReal a, VF_FieldReal b,
     VF_FieldReal tangentA, VF_FieldReal tangentB, VF_FieldReal length,
@@ -249,6 +250,193 @@ VF_FIELD_INLINE float VF_FieldSurfaceRelief(float x, float y, float z, VF_FieldU
     return 1.20f * macro + (0.80f + 0.25f * macro) * (0.40f * mid + 0.60f * fine);
 }
 
+// Scalar version of VoxelNoise::Perlin3D and its owner-world roughness variant.
+// The integer hash and operation order are shared so the standalone field and HLSL
+// evaluate the same authored detail from the same world parameters.
+VF_FIELD_INLINE VF_FieldU32 VF_FieldVoxelHashMix(VF_FieldU32 x)
+{
+    x ^= x >> 16; x *= 0x45d9f3bu;
+    x ^= x >> 16; x *= 0x45d9f3bu;
+    x ^= x >> 16;
+    return x;
+}
+
+VF_FIELD_INLINE float VF_FieldVoxelSeedOffset(VF_FieldU32 seed,float siteKey)
+{
+    const VF_FieldU32 site=VF_FIELD_CAST_U32(siteKey*100.0f+0.5f);
+    return VF_FIELD_CAST_F32(VF_FieldVoxelHashMix(seed^(site*2654435761u))&0x3FFFu);
+}
+
+VF_FIELD_INLINE VF_FieldU32 VF_FieldPerlinHash(VF_FieldI32 x,VF_FieldI32 y,VF_FieldI32 z)
+{
+    VF_FieldU32 h=VF_FIELD_CAST_U32(x)*0x9E3779B1u
+        ^VF_FIELD_CAST_U32(y)*0x85EBCA77u
+        ^VF_FIELD_CAST_U32(z)*0xC2B2AE3Du;
+    h^=h>>15;h*=0x2C1B3C6Du;
+    h^=h>>12;h*=0x297A2D39u;
+    h^=h>>15;
+    return h;
+}
+
+VF_FIELD_INLINE float VF_FieldPerlinFade(float t)
+{
+    const float inner=t*(t*6.0f-15.0f)+10.0f;
+    const float t3=t*t*t;
+    return t3*inner;
+}
+
+VF_FIELD_INLINE float VF_FieldPerlinGrad(VF_FieldU32 hash,float x,float y,float z)
+{
+    const VF_FieldU32 h=hash&15u;
+    const float u=(h&8u)==0u?x:y;
+    float v;
+    if((h&12u)==0u)v=y;
+    else if((h&13u)==12u)v=x;
+    else v=z;
+    const float ru=(h&1u)==0u?u:-u;
+    const float rv=(h&2u)==0u?v:-v;
+    return ru+rv;
+}
+
+VF_FIELD_INLINE float VF_FieldPerlinLerp(float a,float b,float t)
+{ return a+t*(b-a); }
+
+VF_FIELD_INLINE float VF_FieldPerlin3D(float x,float y,float z)
+{
+    const float xf=VF_FIELD_FLOOR(x),yf=VF_FIELD_FLOOR(y),zf=VF_FIELD_FLOOR(z);
+    const VF_FieldI32 X=VF_FIELD_CAST_I32(xf),Y=VF_FIELD_CAST_I32(yf),Z=VF_FIELD_CAST_I32(zf);
+    const float fx=x-xf,fy=y-yf,fz=z-zf;
+    const float su=VF_FieldPerlinFade(fx),sv=VF_FieldPerlinFade(fy),sw=VF_FieldPerlinFade(fz);
+    const VF_FieldU32 h000=VF_FieldPerlinHash(X,Y,Z),h100=VF_FieldPerlinHash(X+1,Y,Z);
+    const VF_FieldU32 h010=VF_FieldPerlinHash(X,Y+1,Z),h110=VF_FieldPerlinHash(X+1,Y+1,Z);
+    const VF_FieldU32 h001=VF_FieldPerlinHash(X,Y,Z+1),h101=VF_FieldPerlinHash(X+1,Y,Z+1);
+    const VF_FieldU32 h011=VF_FieldPerlinHash(X,Y+1,Z+1),h111=VF_FieldPerlinHash(X+1,Y+1,Z+1);
+    const float fx1=fx-1.0f,fy1=fy-1.0f,fz1=fz-1.0f;
+    const float n000=VF_FieldPerlinGrad(h000,fx,fy,fz);
+    const float n100=VF_FieldPerlinGrad(h100,fx1,fy,fz);
+    const float n010=VF_FieldPerlinGrad(h010,fx,fy1,fz);
+    const float n110=VF_FieldPerlinGrad(h110,fx1,fy1,fz);
+    const float n001=VF_FieldPerlinGrad(h001,fx,fy,fz1);
+    const float n101=VF_FieldPerlinGrad(h101,fx1,fy,fz1);
+    const float n011=VF_FieldPerlinGrad(h011,fx,fy1,fz1);
+    const float n111=VF_FieldPerlinGrad(h111,fx1,fy1,fz1);
+    const float x00=VF_FieldPerlinLerp(n000,n100,su);
+    const float x10=VF_FieldPerlinLerp(n010,n110,su);
+    const float x01=VF_FieldPerlinLerp(n001,n101,su);
+    const float x11=VF_FieldPerlinLerp(n011,n111,su);
+    const float y0=VF_FieldPerlinLerp(x00,x10,sv);
+    const float y1=VF_FieldPerlinLerp(x01,x11,sv);
+    return VF_FieldPerlinLerp(y0,y1,sw);
+}
+
+VF_FIELD_INLINE float VF_FieldOwnerFractal(float x,float y,float z,VF_FieldI32 octaves,
+    VF_FieldI32 mode)
+{
+    float total=0.0f,maxValue=0.0f,frequency=1.0f,amplitude=1.0f,weight=1.0f;
+    for(VF_FieldI32 i=0;i<octaves;++i)
+    {
+        const float noise=VF_FieldPerlin3D(x*frequency,y*frequency,z*frequency);
+        if(mode==0)total+=noise*amplitude;
+        else
+        {
+            float ridge=noise*1.25f;
+            ridge=1.0f-VF_FIELD_ABS(ridge);
+            ridge=ridge*ridge;
+            ridge*=weight;
+            weight=VF_FIELD_CLAMP(ridge*2.0f,0.0f,1.0f);
+            total+=ridge*amplitude;
+        }
+        maxValue+=amplitude;
+        amplitude*=0.5f;
+        frequency*=2.0f;
+    }
+    const float result=total/maxValue;
+    return mode==0?result:result*2.0f-1.0f;
+}
+
+VF_FIELD_INLINE float VF_FieldOwnerCellular(float x,float y,float z)
+{
+    const VF_FieldI32 cellX=VF_FIELD_CAST_I32(VF_FIELD_FLOOR(x));
+    const VF_FieldI32 cellY=VF_FIELD_CAST_I32(VF_FIELD_FLOOR(y));
+    const VF_FieldI32 cellZ=VF_FIELD_CAST_I32(VF_FIELD_FLOOR(z));
+    const float fracX=x-VF_FIELD_CAST_F32(cellX),fracY=y-VF_FIELD_CAST_F32(cellY),
+        fracZ=z-VF_FIELD_CAST_F32(cellZ);
+    float first=3.402823466e+38f,second=3.402823466e+38f;
+    for(VF_FieldI32 dz=-1;dz<=1;++dz)
+    for(VF_FieldI32 dy=-1;dy<=1;++dy)
+    for(VF_FieldI32 dx=-1;dx<=1;++dx)
+    {
+        const VF_FieldI32 nx=cellX+dx,ny=cellY+dy,nz=cellZ+dz;
+        const VF_FieldU32 h=VF_FieldVoxelHashMix(
+            VF_FIELD_CAST_U32(nx+0x7FFFFFFF)
+            ^VF_FieldVoxelHashMix(VF_FIELD_CAST_U32(ny+0x7FFFFFFF)*2654435761u)
+            ^VF_FieldVoxelHashMix(VF_FIELD_CAST_U32(nz+0x7FFFFFFF)*374761393u));
+        const float fpx=VF_FIELD_CAST_F32(dx)
+            +VF_FIELD_CAST_F32(h&0xFFFFFFu)*(1.0f/16777216.0f)-fracX;
+        const VF_FieldU32 hy=VF_FieldVoxelHashMix(h^0x12345678u);
+        const VF_FieldU32 hz=VF_FieldVoxelHashMix(h^0x9ABCDEF0u);
+        const float fpy=VF_FIELD_CAST_F32(dy)
+            +VF_FIELD_CAST_F32(hy&0xFFFFFFu)*(1.0f/16777216.0f)-fracY;
+        const float fpz=VF_FIELD_CAST_F32(dz)
+            +VF_FIELD_CAST_F32(hz&0xFFFFFFu)*(1.0f/16777216.0f)-fracZ;
+        const float distanceSquared=fpx*fpx+fpy*fpy+fpz*fpz;
+        if(distanceSquared<first){second=first;first=distanceSquared;}
+        else if(distanceSquared<second)second=distanceSquared;
+    }
+    return (VF_FIELD_SQRT(second)-VF_FIELD_SQRT(first))*2.0f-1.0f;
+}
+
+VF_FIELD_INLINE float VF_FieldOwnerRoughnessCore(float x,float y,float z,
+    VF_FieldU32 seed,float surfaceRoughness,float roughnessFrequency,float verticalScale,
+    float domainWarpStrength,float domainWarpFrequency,VF_FieldU32 noiseType)
+{
+    if(!(roughnessFrequency>0.0f))return 0.0f;
+    const float effectiveZ=verticalScale>0.0f&&verticalScale!=1.0f?z/verticalScale:z;
+    const float mainX=x*roughnessFrequency+VF_FieldVoxelSeedOffset(seed,11.3f);
+    const float mainY=y*roughnessFrequency+VF_FieldVoxelSeedOffset(seed,13.7f);
+    const float mainZ=effectiveZ*roughnessFrequency+VF_FieldVoxelSeedOffset(seed,17.1f);
+    float fineX=x*roughnessFrequency*3.0f+VF_FieldVoxelSeedOffset(seed,19.1f)+2000.0f;
+    float fineY=y*roughnessFrequency*3.0f+VF_FieldVoxelSeedOffset(seed,23.7f)+2500.0f;
+    float fineZ=effectiveZ*roughnessFrequency*3.0f+VF_FieldVoxelSeedOffset(seed,29.3f)+3000.0f;
+    float mx=mainX,my=mainY,mz=mainZ;
+    if(domainWarpStrength>0.0f)
+    {
+        const float wf=domainWarpFrequency,ws=domainWarpStrength;
+        float wx=VF_FieldPerlin3D(x*wf+VF_FieldVoxelSeedOffset(seed,5.2f),
+            y*wf+VF_FieldVoxelSeedOffset(seed,1.3f),
+            effectiveZ*wf+VF_FieldVoxelSeedOffset(seed,9.7f));
+        float wy=VF_FieldPerlin3D(x*wf+100.0f+VF_FieldVoxelSeedOffset(seed,7.7f),
+            y*wf+200.0f+VF_FieldVoxelSeedOffset(seed,3.1f),effectiveZ*wf+300.0f);
+        float wz=VF_FieldPerlin3D(x*wf+400.0f,y*wf+500.0f+VF_FieldVoxelSeedOffset(seed,11.9f),
+            effectiveZ*wf+600.0f+VF_FieldVoxelSeedOffset(seed,13.3f));
+        wx*=1.25f*ws;wy*=1.25f*ws;wz*=1.25f*ws;
+        mx+=wx;my+=wy;mz+=wz;
+        fineX+=wx;fineY+=wy;fineZ+=wz;
+    }
+    float rough=0.0f,fine=0.0f;
+    if(noiseType==1u)
+    {
+        rough=VF_FieldOwnerFractal(mx,my,mz,3,1);
+        fine=VF_FieldOwnerFractal(fineX,fineY,fineZ,2,1);
+    }
+    else if(noiseType==2u)
+    {
+        rough=(VF_FieldOwnerFractal(mx,my,mz,3,0)+VF_FieldOwnerFractal(mx,my,mz,3,1))*0.5f;
+        fine=(VF_FieldOwnerFractal(fineX,fineY,fineZ,2,0)+VF_FieldOwnerFractal(fineX,fineY,fineZ,2,1))*0.5f;
+    }
+    else if(noiseType==3u)
+    {
+        rough=VF_FieldOwnerCellular(mx,my,mz);
+        fine=VF_FieldOwnerCellular(fineX,fineY,fineZ);
+    }
+    else
+    {
+        rough=VF_FieldOwnerFractal(mx,my,mz,3,0);
+        fine=VF_FieldOwnerFractal(fineX,fineY,fineZ,2,0);
+    }
+    return (rough+fine*0.4f)*1.25f*surfaceRoughness;
+}
+
 VF_FIELD_INLINE float VF_FieldSmoothUnit(float value)
 {
 #if defined(VF_FIELD_HLSL)

Standalone/Core/verification/RunStep6CpuHash.ps1

@@ -0,0 +1,51 @@
+[CmdletBinding()]
+param()
+
+$ErrorActionPreference = 'Stop'
+$HostProcess = Get-Process -Id $PID
+$HostProcess.PriorityClass = 'BelowNormal'
+$HostProcess.ProcessorAffinity = [IntPtr]0xFF00
+$Standalone = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..'))
+$Executable = Join-Path $Standalone 'bin\VoxelForgeGpuFieldStep6NearProbe.exe'
+$Stdout = Join-Path $PSScriptRoot 'step6_cpu_hash_fresh_w8.stdout.log'
+$Stderr = Join-Path $PSScriptRoot 'step6_cpu_hash_fresh_w8.stderr.log'
+$Arguments = @('--stream-benchmark','--shape-world','--force-generate','--stream-hash','--workers=8','--clip-radius=3')
+if (-not (Test-Path -LiteralPath $Executable -PathType Leaf)) { throw "Executable not found: $Executable" }
+
+for ($Attempt = 0; $Attempt -lt 2; ++$Attempt) {
+    if ($Attempt -eq 1) {
+        Write-Output '0xC0000142 encountered; waiting 60 seconds before the permitted retry.'
+        Start-Sleep -Seconds 60
+    }
+    $Child = Start-Process -FilePath $Executable -ArgumentList $Arguments -WorkingDirectory $Standalone `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $Stdout -RedirectStandardError $Stderr
+    $PeakWorkingSet = 0L
+    $PeakPrivate = 0L
+    $OverLimit = $false
+    while (-not $Child.HasExited) {
+        $Child.Refresh()
+        if ($Child.HasExited) { break }
+        $Child.PriorityClass = 'BelowNormal'
+        $Child.ProcessorAffinity = [IntPtr]0xFF00
+        $PeakWorkingSet = [Math]::Max($PeakWorkingSet,$Child.WorkingSet64)
+        $PeakPrivate = [Math]::Max($PeakPrivate,$Child.PrivateMemorySize64)
+        if ($Child.PrivateMemorySize64 -ge 8GB) {
+            $OverLimit = $true
+            $Child.Kill()
+            $Child.WaitForExit()
+            break
+        }
+        Start-Sleep -Milliseconds 100
+    }
+    $Child.WaitForExit()
+    $Child.Refresh()
+    if ($Child.ExitCode -eq -1073741502 -and $Attempt -eq 0) { continue }
+    if ($Child.ExitCode -eq -1073741502) { throw '0xC0000142 repeated after the one retry; stopping.' }
+    if ($OverLimit) { throw 'CPU hash run stopped at the 8 GB process memory limit.' }
+    $Log = Get-Content -LiteralPath $Stdout -Raw
+    if ($Child.ExitCode -ne 0 -or -not $Log.Contains('[StreamFullClipmap]')) {
+        throw "Hash run failed; exit=$($Child.ExitCode) stdout=$Stdout stderr=$Stderr"
+    }
+    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($PeakPrivate/1GB,2)) exit_code=$($Child.ExitCode) stdout=$Stdout stderr=$Stderr"
+    exit 0
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h

@@ -13,6 +13,7 @@ struct ID3D11Query;
 
 namespace VoxelForgeGpuFarViewer
 {
+    inline constexpr std::uint32_t PooledOutputVertexCapacityPerTile=19456;
     struct CpuVertex { float position[3]{}, normal[3]{}; };
     struct CpuMesh { std::vector<CpuVertex> vertices; std::vector<std::uint32_t> indices; };
     struct GpuTimingQueries

Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp

@@ -128,17 +128,21 @@ namespace
         std::uint32_t cells, sampleDim, cellCount, groupCount;
         std::int32_t originX, originY, originZ, stepVoxels;
         float voxelSizeCm;
-        std::uint32_t outputBaseVertex=0, outputArgsByteOffset=0, pad=0;
+        std::uint32_t outputBaseVertex=0, outputArgsByteOffset=0, nearFieldMode=0;
+        std::uint32_t outputVertexCapacity=0, padding[3]{};
     };
-    static_assert(sizeof(TileConstants) == 48);
+    static_assert(sizeof(TileConstants) == 64);
+    static_assert(offsetof(TileConstants, outputVertexCapacity) == 48);
 
     struct FarFieldConstants
     {
         float strataLo, strataHi, strataBlend;
         std::uint32_t sampleCount, evalCounterEnabled, spatialBinMask, spatialBinSize,
             spatialBinCapacity;
+        std::uint32_t nearDetailEnabled, fieldSeed, detailLayerCount, pad;
+        std::int32_t detailMinVoxel[4], detailMaxVoxel[4];
     };
-    static_assert(sizeof(FarFieldConstants) == 32);
+    static_assert(sizeof(FarFieldConstants) == 80);
     static_assert(sizeof(FarFieldPrimitiveGpu) == 224);
     static_assert(sizeof(FarFieldSpatialBin) == 24);
     static_assert(sizeof(FarFieldStratumBoundary) == 8);
@@ -475,6 +479,25 @@ namespace
                 detail.assign(static_cast<const char*>(errors->GetBufferPointer()), errors->GetBufferSize());
             throw std::runtime_error(std::string("D3DCompile ") + entry + " failed: " + detail);
         }
+        if(std::string_view(entry)=="EvaluateFarFieldCS")
+        {
+            ComPtr<ID3D11ShaderReflection> reflection;
+            if(SUCCEEDED(D3DReflect(bytecode->GetBufferPointer(),bytecode->GetBufferSize(),
+                __uuidof(ID3D11ShaderReflection),reinterpret_cast<void**>(reflection.GetAddressOf()))))
+            {
+                ID3D11ShaderReflectionConstantBuffer* cbuffer=
+                    reflection->GetConstantBufferByName("FarFieldConstants");
+                ID3D11ShaderReflectionVariable* variable=cbuffer
+                    ?cbuffer->GetVariableByName("FarNearDetailEnabled"):nullptr;
+                D3D11_SHADER_VARIABLE_DESC variableDesc{};
+                D3D11_SHADER_BUFFER_DESC bufferDesc{};
+                if(cbuffer&&variable&&SUCCEEDED(cbuffer->GetDesc(&bufferDesc))
+                    &&SUCCEEDED(variable->GetDesc(&variableDesc)))
+                    std::printf("[NearShaderLayoutProbe] cbuffer_size=%u near_detail_offset=%u cpp_offset=%zu\n",
+                        bufferDesc.Size,variableDesc.StartOffset,
+                        offsetof(FarFieldConstants,nearDetailEnabled));
+            }
+        }
         return bytecode;
     }
 
@@ -496,36 +519,38 @@ namespace
         struct FarScratch
         {
             std::uint32_t sampleCount=0,cellCount=0,groupCount=0;
-            ComPtr<ID3D11Buffer> signs,tri,density,midpointIds,midpointRecords;
+            ComPtr<ID3D11Buffer> signs,tri,density,nearNormalDensity,midpointIds,midpointRecords;
             ComPtr<ID3D11Buffer> cases,localOffsets,groupCounts,groupOffsets;
             ComPtr<ID3D11Buffer> tileConstants,farConstants;
             ComPtr<ID3D11Buffer> connectorConstants,connectorRings;
-            ComPtr<ID3D11ShaderResourceView> signSrv,triSrv,densitySrv,midpointIdSrv,midpointRecordSrv;
+            ComPtr<ID3D11ShaderResourceView> signSrv,triSrv,densitySrv,nearNormalDensitySrv,midpointIdSrv,midpointRecordSrv;
             ComPtr<ID3D11ShaderResourceView> caseSrv,localOffsetSrv,groupCountSrv,groupOffsetSrv;
             ComPtr<ID3D11ShaderResourceView> connectorRingSrv;
-            ComPtr<ID3D11UnorderedAccessView> densityUav,caseUav,localOffsetUav,groupCountUav,groupOffsetUav;
+            ComPtr<ID3D11UnorderedAccessView> densityUav,nearNormalDensityUav,caseUav,localOffsetUav,groupCountUav,groupOffsetUav;
             ComPtr<ID3D11UnorderedAccessView> connectorRingUav;
-            DynamicInput shapes,ranges,ids,protectedAir,protectedSolid;
+            DynamicInput shapes,ranges,ids,protectedAir,protectedSolid,nearPackedProxy;
             DynamicInput connectors;
             std::uint32_t connectorRingCapacity=0;
         } farScratch;
         struct FarWorldResources
         {
             std::shared_ptr<const FarFieldWorldInput> owner;
-            DynamicInput shapes,bins,shapeIds;
+            DynamicInput shapes,bins,shapeIds,detailLayers;
         } farWorldResources;
         ComPtr<ID3D11Device> device;
         ComPtr<ID3D11DeviceContext> context;
-        ComPtr<ID3D11DeviceContext> immediateContext,deferredContext;
+            ComPtr<ID3D11DeviceContext> immediateContext,deferredContext;
         ComPtr<ID3D11ComputeShader> fieldEval, classify, fieldClassify, scan, emit, noiseProbe;
         ComPtr<ID3D11ComputeShader> connectorRings,connectorPanels;
         ComPtr<IDXGIAdapter3> memoryAdapter;
         std::string name;
         UINT64 timestampFrequency = 0;
         std::vector<FarOutputChunk> farOutputChunks;
-        std::uint32_t farOutputCapacitySlots=0,farOutputCells=0;
+        std::uint32_t farOutputCapacitySlots=0;
         std::uint32_t farOutputSlotsPerChunk=128,farMaxVerticesPerTile=0;
-        static constexpr std::uint32_t FarPanelVertexCapacity=32768;
+        static constexpr std::uint32_t FarOutputVertexCapacityPerTile=
+            VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile;
+        static constexpr std::uint32_t FarPanelVertexCapacity=28160;
 
         void CompileShaders(const std::filesystem::path& shaderPath)
         {
@@ -731,13 +756,9 @@ namespace
         void EnsureFarOutputPool(std::uint32_t tileSlots,std::uint32_t cells)
         {
             if(tileSlots==0||cells==0)throw std::runtime_error("far output pool dimensions must be non-zero");
-            if(tileSlots<=farOutputCapacitySlots&&cells==farOutputCells)return;
+            if(tileSlots<=farOutputCapacitySlots&&!farOutputChunks.empty())return;
-            const std::uint64_t maxVertices64=static_cast<std::uint64_t>(cells)*cells*cells*15ull;
+            farMaxVerticesPerTile=FarOutputVertexCapacityPerTile;
-            if(maxVertices64>std::numeric_limits<std::uint32_t>::max())
-                throw std::runtime_error("far output tile vertex capacity exceeds 32-bit limits");
-            farMaxVerticesPerTile=static_cast<std::uint32_t>(maxVertices64);
             farOutputCapacitySlots=std::max(tileSlots,farOutputCapacitySlots);
-            farOutputCells=cells;
             const std::uint32_t chunkCount=(farOutputCapacitySlots+farOutputSlotsPerChunk-1)/farOutputSlotsPerChunk;
             std::vector<FarOutputChunk> chunks;
             chunks.reserve(chunkCount);
@@ -819,6 +840,8 @@ namespace
                 "Map(world far spatial bins)");
             UpdateDynamicInput(farWorldResources.shapeIds,input->spatialBinShapeIds.data(),
                 input->spatialBinShapeIds.size(),"Map(world far bin shape IDs)");
+            UpdateDynamicInput(farWorldResources.detailLayers,input->detailLayers.data(),
+                input->detailLayers.size(),"Map(world near detail layers)");
             farWorldResources.owner=input;
         }
 
@@ -837,6 +860,8 @@ namespace
                 for(std::size_t s=0;s<16;++s)triData[c*16+s]=MCTriTable[c][s];
             scratch.tri=Structured(triData.data(),triData.size(),D3D11_BIND_SHADER_RESOURCE);
             scratch.density=Structured<float>(nullptr,samples,D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
+            scratch.nearNormalDensity=Structured<float>(nullptr,samples,
+                D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
             std::vector<std::uint32_t> invalidIds(cells,InvalidIndex);
             scratch.midpointIds=Structured(invalidIds.data(),invalidIds.size(),D3D11_BIND_SHADER_RESOURCE);
             const MidpointRecord zeroRecord{};
@@ -848,6 +873,7 @@ namespace
             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);
+            scratch.nearNormalDensitySrv=Srv<float>(scratch.nearNormalDensity.Get(),samples);
             scratch.midpointIdSrv=Srv<std::uint32_t>(scratch.midpointIds.Get(),cells);
             scratch.midpointRecordSrv=Srv<MidpointRecord>(scratch.midpointRecords.Get(),1);
             scratch.caseSrv=Srv<std::uint32_t>(scratch.cases.Get(),cells);
@@ -855,6 +881,7 @@ namespace
             scratch.groupCountSrv=Srv<std::uint32_t>(scratch.groupCounts.Get(),groups);
             scratch.groupOffsetSrv=Srv<std::uint32_t>(scratch.groupOffsets.Get(),groups);
             scratch.densityUav=Uav(scratch.density.Get(),samples);
+            scratch.nearNormalDensityUav=Uav(scratch.nearNormalDensity.Get(),samples);
             scratch.caseUav=Uav(scratch.cases.Get(),cells);
             scratch.localOffsetUav=Uav(scratch.localOffsets.Get(),cells);
             scratch.groupCountUav=Uav(scratch.groupCounts.Get(),groups);
@@ -949,7 +976,7 @@ namespace
             FarOutputChunk* pooledChunk=nullptr;
             if(usePooledFarOutput)
             {
-                if(!farField||farOutputCells!=tile.cells||outputSlot>=farOutputCapacitySlots
+                if(!farField||outputSlot>=farOutputCapacitySlots
                     ||farOutputChunks.empty())
                     throw std::runtime_error("far output pool was not reserved for this tile slot");
                 const std::uint32_t chunkIndex=outputSlot/farOutputSlotsPerChunk;
@@ -977,6 +1004,7 @@ namespace
             if (farField)
                 result.d3dBufferBytes += farField->protectedAir.size() * sizeof(std::uint32_t)
                     + farField->protectedSolid.size() * sizeof(std::uint32_t)
+                    + farField->nearPackedProxy.size() * sizeof(std::uint32_t)
                     + farField->connectorPanels.size() * sizeof(FarFieldConnectorPanel)
                     + tile.density.size() * sizeof(float)
                     + sizeof(FarFieldConstants);
@@ -985,14 +1013,14 @@ namespace
                     + std::max<std::uint32_t>(1, tile.expectedVertexCount) * sizeof(GpuVertex)
                     + std::max<std::uint32_t>(1, tile.expectedVertexCount) * sizeof(std::uint32_t)
                     + 5 * sizeof(std::uint32_t);
-            ComPtr<ID3D11Buffer> signs,triBuffer,density,midpointIds,midpointRecords;
+            ComPtr<ID3D11Buffer> signs,triBuffer,density,nearNormalDensity,midpointIds,midpointRecords;
             ComPtr<ID3D11Buffer> cases,localOffsets,groupCounts,groupOffsets;
-            ComPtr<ID3D11ShaderResourceView> signSrv,triSrv,densitySrv,farShapeSrv,farRangeSrv,farIdSrv;
+            ComPtr<ID3D11ShaderResourceView> signSrv,triSrv,densitySrv,nearNormalDensitySrv,farShapeSrv,farRangeSrv,farIdSrv;
-            ComPtr<ID3D11ShaderResourceView> farProtectedSrv,farProtectedSolidSrv;
+            ComPtr<ID3D11ShaderResourceView> farProtectedSrv,farProtectedSolidSrv,farNearProxySrv;
             ComPtr<ID3D11ShaderResourceView> farConnectorSrv,farConnectorRingSrv;
             ComPtr<ID3D11ShaderResourceView> midpointIdSrv,midpointRecordSrv,caseSrv,localOffsetSrv;
             ComPtr<ID3D11ShaderResourceView> groupCountSrv,groupOffsetSrv;
-            ComPtr<ID3D11UnorderedAccessView> caseUav,localOffsetUav,groupCountUav,groupOffsetUav,densityUav;
+            ComPtr<ID3D11UnorderedAccessView> caseUav,localOffsetUav,groupCountUav,groupOffsetUav,densityUav,nearNormalDensityUav;
             ComPtr<ID3D11UnorderedAccessView> farConnectorRingUav;
             ComPtr<ID3D11Buffer> farConnectorConstantsBuffer;
             ComPtr<ID3D11Buffer> farConnectorRingBuffer;
@@ -1010,14 +1038,17 @@ namespace
                 EnsureFarScratch(tile);
                 FarScratch& scratch=farScratch;
                 signs=scratch.signs;triBuffer=scratch.tri;density=scratch.density;
+                nearNormalDensity=scratch.nearNormalDensity;
                 midpointIds=scratch.midpointIds;midpointRecords=scratch.midpointRecords;
                 cases=scratch.cases;localOffsets=scratch.localOffsets;
                 groupCounts=scratch.groupCounts;groupOffsets=scratch.groupOffsets;
                 signSrv=scratch.signSrv;triSrv=scratch.triSrv;densitySrv=scratch.densitySrv;
+                nearNormalDensitySrv=scratch.nearNormalDensitySrv;
                 midpointIdSrv=scratch.midpointIdSrv;midpointRecordSrv=scratch.midpointRecordSrv;
                 caseSrv=scratch.caseSrv;localOffsetSrv=scratch.localOffsetSrv;
                 groupCountSrv=scratch.groupCountSrv;groupOffsetSrv=scratch.groupOffsetSrv;
-                densityUav=scratch.densityUav;caseUav=scratch.caseUav;
+                densityUav=scratch.densityUav;nearNormalDensityUav=scratch.nearNormalDensityUav;
+                caseUav=scratch.caseUav;
                 localOffsetUav=scratch.localOffsetUav;groupCountUav=scratch.groupCountUav;
                 groupOffsetUav=scratch.groupOffsetUav;
                 farShapeSrv=farWorldResources.shapes.srv;
@@ -1027,6 +1058,11 @@ namespace
                     farField->protectedAir.data(),farField->protectedAir.size(),"Map(far protected air)");
                 farProtectedSolidSrv=UpdateDynamicInput(scratch.protectedSolid,
                     farField->protectedSolid.data(),farField->protectedSolid.size(),"Map(far protected solid)");
+                const std::uint32_t noPackedProxy=std::numeric_limits<std::uint32_t>::max();
+                farNearProxySrv=UpdateDynamicInput(scratch.nearPackedProxy,
+                    farField->nearPackedProxy.empty()?&noPackedProxy:farField->nearPackedProxy.data(),
+                    farField->nearPackedProxy.empty()?1:farField->nearPackedProxy.size(),
+                    "Map(far near packed proxies)");
                 farConnectorSrv=UpdateDynamicInput(scratch.connectors,
                     farField->connectorPanels.empty()?nullptr:farField->connectorPanels.data(),
                     farField->connectorPanels.size(),"Map(far connector panels)");
@@ -1041,6 +1077,9 @@ namespace
                 triBuffer=Structured(tri.data(),triCount,D3D11_BIND_SHADER_RESOURCE);
                 density=Structured(farField?static_cast<const float*>(nullptr):tile.density.data(),
                     tile.density.size(),D3D11_BIND_SHADER_RESOURCE|(farField?D3D11_BIND_UNORDERED_ACCESS:0));
+                if(farField)
+                    nearNormalDensity=Structured<float>(nullptr,tile.density.size(),
+                        D3D11_BIND_SHADER_RESOURCE|D3D11_BIND_UNORDERED_ACCESS);
                 midpointIds=Structured(tile.midpointIds.data(),tile.midpointIds.size(),D3D11_BIND_SHADER_RESOURCE);
                 const MidpointRecord zeroRecord{};
                 midpointRecords=Structured(tile.midpointRecords.empty()?&zeroRecord:tile.midpointRecords.data(),
@@ -1052,6 +1091,8 @@ namespace
                 signSrv=Srv<std::uint32_t>(signs.Get(),tile.signs.size());
                 triSrv=Srv<std::int32_t>(triBuffer.Get(),triCount);
                 densitySrv=Srv<float>(density.Get(),tile.density.size());
+                if(farField)
+                    nearNormalDensitySrv=Srv<float>(nearNormalDensity.Get(),tile.density.size());
                 midpointIdSrv=Srv<std::uint32_t>(midpointIds.Get(),tile.midpointIds.size());
                 midpointRecordSrv=Srv<MidpointRecord>(midpointRecords.Get(),std::max<std::size_t>(1,tile.midpointRecords.size()));
                 caseSrv=Srv<std::uint32_t>(cases.Get(),cellCount);
@@ -1061,6 +1102,8 @@ namespace
                 caseUav=Uav(cases.Get(),cellCount);localOffsetUav=Uav(localOffsets.Get(),cellCount);
                 groupCountUav=Uav(groupCounts.Get(),groupCount);groupOffsetUav=Uav(groupOffsets.Get(),groupCount);
                 densityUav=farField?Uav(density.Get(),tile.density.size()):ComPtr<ID3D11UnorderedAccessView>{};
+                nearNormalDensityUav=farField?Uav(nearNormalDensity.Get(),tile.density.size()):
+                    ComPtr<ID3D11UnorderedAccessView>{};
                 if(farField)
                 {
                     const std::uint32_t zeroWord=0;
@@ -1068,11 +1111,18 @@ namespace
                         std::max<std::size_t>(1,farField->protectedAir.size()),D3D11_BIND_SHADER_RESOURCE);
                     auto farProtectedSolid=Structured(farField->protectedSolid.empty()?&zeroWord:farField->protectedSolid.data(),
                         std::max<std::size_t>(1,farField->protectedSolid.size()),D3D11_BIND_SHADER_RESOURCE);
+                    const std::uint32_t noPackedProxy=std::numeric_limits<std::uint32_t>::max();
+                    auto farNearProxy=Structured(farField->nearPackedProxy.empty()?&noPackedProxy:
+                        farField->nearPackedProxy.data(),std::max<std::size_t>(1,
+                        farField->nearPackedProxy.empty()?1:farField->nearPackedProxy.size()),
+                        D3D11_BIND_SHADER_RESOURCE);
                     farShapeSrv=farWorldResources.shapes.srv;
                     farRangeSrv=farWorldResources.bins.srv;
                     farIdSrv=farWorldResources.shapeIds.srv;
                     farProtectedSrv=Srv<std::uint32_t>(farProtected.Get(),std::max<std::size_t>(1,farField->protectedAir.size()));
                     farProtectedSolidSrv=Srv<std::uint32_t>(farProtectedSolid.Get(),std::max<std::size_t>(1,farField->protectedSolid.size()));
+                    farNearProxySrv=Srv<std::uint32_t>(farNearProxy.Get(),std::max<std::size_t>(1,
+                        farField->nearPackedProxy.empty()?1:farField->nearPackedProxy.size()));
                     const FarFieldConnectorPanel zeroPanel{};
                     auto farConnectors=Structured(farField->connectorPanels.empty()?&zeroPanel:
                         farField->connectorPanels.data(),std::max<std::size_t>(1,
@@ -1125,18 +1175,26 @@ namespace
             cbDesc.BindFlags = D3D11_BIND_CONSTANT_BUFFER;
             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,0};
+                25.0f,outputBaseVertex,outputArgsByteOffset,
+                farField&&farField->nearDetailEnabled?1u:0u,
+                pooledChunk?farMaxVerticesPerTile:static_cast<std::uint32_t>(maxVertices),{0,0,0}};
             result.outputBaseVertex=outputBaseVertex;
             result.outputArgsByteOffset=outputArgsByteOffset;
             ComPtr<ID3D11Buffer> cb;
             ComPtr<ID3D11Buffer> farCb;
             if (farField)
             {
-                const FarFieldConstants farConstants{farField->strataLo,farField->strataHi,
+                FarFieldConstants farConstants{farField->strataLo,farField->strataHi,
                     farField->strataBlend,static_cast<std::uint32_t>(tile.density.size()),
                     profileGpuKernels?1u:0u,farField->worldInput->spatialBinMask,
                     farField->worldInput->spatialBinSize,
-                    static_cast<std::uint32_t>(farField->worldInput->spatialBins.size())};
+                    static_cast<std::uint32_t>(farField->worldInput->spatialBins.size()),
+                    farField->nearDetailEnabled?1u:0u,farField->worldInput->fieldSeed,
+                    static_cast<std::uint32_t>(farField->worldInput->detailLayers.size()),0u};
+                std::copy(std::begin(farField->worldInput->detailMinVoxel),
+                    std::end(farField->worldInput->detailMinVoxel),farConstants.detailMinVoxel);
+                std::copy(std::begin(farField->worldInput->detailMaxVoxel),
+                    std::end(farField->worldInput->detailMaxVoxel),farConstants.detailMaxVoxel);
                 const FarConnectorConstants connectorConstants{
                     static_cast<std::uint32_t>(farField->connectorPanels.size()),
                     farField->connectorRingCount,farField->connectorVertexCount,connectorPanelCapacity,
@@ -1223,19 +1281,22 @@ namespace
                 context->CSSetConstantBuffers(1,1,farCb.GetAddressOf());
                 context->CSSetConstantBuffers(2,1,farConnectorConstantsBuffer.GetAddressOf());
                 context->CSSetShader(fieldEval.Get(),nullptr,0);
-                ID3D11ShaderResourceView* fieldSrvs[15]{};
+                ID3D11ShaderResourceView* fieldSrvs[18]{};
                 fieldSrvs[10]=farShapeSrv.Get();fieldSrvs[11]=farRangeSrv.Get();
                 fieldSrvs[12]=farIdSrv.Get();fieldSrvs[13]=farProtectedSrv.Get();
                 fieldSrvs[14]=farProtectedSolidSrv.Get();
-                context->CSSetShaderResources(0,15,fieldSrvs);
+                fieldSrvs[16]=farWorldResources.detailLayers.srv.Get();
-                ID3D11UnorderedAccessView* fieldUavs[8]{};fieldUavs[6]=farEvalCounterUav.Get();
+                fieldSrvs[17]=farNearProxySrv.Get();
+                context->CSSetShaderResources(0,18,fieldSrvs);
+                ID3D11UnorderedAccessView* fieldUavs[8]{};fieldUavs[5]=nearNormalDensityUav.Get();
+                fieldUavs[6]=farEvalCounterUav.Get();
                 fieldUavs[7]=densityUav.Get();
                 context->CSSetUnorderedAccessViews(0,8,fieldUavs,nullptr);
                 context->Dispatch(static_cast<UINT>((tile.density.size()+63)/64),1,1);
                 endKernel(0);
-                ID3D11ShaderResourceView* clearSrvs[16]{};
+                ID3D11ShaderResourceView* clearSrvs[19]{};
                 ID3D11UnorderedAccessView* clearUavs[8]{};
-                context->CSSetShaderResources(0,16,clearSrvs);
+                context->CSSetShaderResources(0,19,clearSrvs);
                 context->CSSetUnorderedAccessViews(0,8,clearUavs,nullptr);
                 if(captureGpuTiming&&!batchTimingMode)context->End(fieldEndQuery.Get());
             }
@@ -1267,16 +1328,18 @@ namespace
 
             beginKernel(3);
             context->CSSetShader(emit.Get(), nullptr, 0);
-            ID3D11ShaderResourceView* emitSrvs[9]{};
+            ID3D11ShaderResourceView* emitSrvs[19]{};
             emitSrvs[1] = triSrv.Get(); emitSrvs[2] = densitySrv.Get(); emitSrvs[3] = caseSrv.Get();
             emitSrvs[4] = localOffsetSrv.Get(); emitSrvs[5] = groupOffsetSrv.Get();
             emitSrvs[6] = midpointIdSrv.Get(); emitSrvs[7] = midpointRecordSrv.Get();
-            context->CSSetShaderResources(0, 9, emitSrvs);
+            if(farField)emitSrvs[18]=nearNormalDensitySrv.Get();
+            context->CSSetShaderResources(0, 19, emitSrvs);
             ID3D11UnorderedAccessView* emitUavs[7]{}; emitUavs[5] = vertexUav.Get(); emitUavs[6] = indexUav.Get();
             context->CSSetUnorderedAccessViews(0, 7, emitUavs, nullptr);
             context->Dispatch(groupCount, 1, 1);
             endKernel(3);
-            context->CSSetShaderResources(0, 16, nullSrvs);
+            ID3D11ShaderResourceView* nullEmitSrvs[19]{};
+            context->CSSetShaderResources(0, 19, nullEmitSrvs);
             context->CSSetUnorderedAccessViews(0, 7, nullUavs, nullptr);
 
             if(farField)
@@ -1742,7 +1805,7 @@ namespace
     }
 
     std::vector<FarSelection> SelectFarTiles(World& world,const WorldConfiguration& config,
-                                             Vec3d camera,bool wholeFar)
+                                             Vec3d camera,bool wholeFar,bool wholeNear)
     {
         std::vector<FarSelection> result;
         const auto desired=GetFullDesired(world,config,camera);
@@ -1762,6 +1825,20 @@ namespace
             std::printf(" mode=all-desired-far-keys lazy-export=1\n");
             return result;
         }
+        if(wholeNear)
+        {
+            std::uint64_t perLod[2]{};
+            for(const TileKey& key:desired)
+            {
+                if(key.lod<0||key.lod>=config.fullResolutionLodCount||key.lod>1)continue;
+                result.push_back({key,{},0.0});
+                ++perLod[key.lod];
+            }
+            std::printf("[WholeNearSelection] desired_near_tiles=%zu lod0=%llu lod1=%llu mode=all-desired-near-keys detail=coherent\n",
+                result.size(),static_cast<unsigned long long>(perLod[0]),
+                static_cast<unsigned long long>(perLod[1]));
+            return result;
+        }
         for(int32_t lod=std::max(config.fullResolutionLodCount,2);lod<=lastLod;++lod)
         {
             std::vector<TileKey> candidates;
@@ -1814,7 +1891,7 @@ namespace
         }
     }
 
-    bool ReportTile(const TileData& tile, const TileMesh& cpuMesh,
+    bool ReportTile(World& world,const TileData& tile, const TileMesh& cpuMesh,
                     const GpuResult& hwFirst, const GpuResult& hwSecond,
                     const GpuResult& warp,const FarFieldTileInput& farInput)
     {
@@ -1868,6 +1945,65 @@ namespace
         const double cpuWarpMean=cpuWarpDirectedCount
             ?(cpuToWarp.meanCm*cpuToWarp.count+warpToCpu.meanCm*warpToCpu.count)/cpuWarpDirectedCount:0.0;
         const float cpuWarpMaximum=std::max(cpuToWarp.maxCm,warpToCpu.maxCm);
+
+        // Connector repair panels are a separate GPU output appended after the ordinary
+        // marching-cubes vertices. Comparing those panels only with CPU surface vertices
+        // conflates two different meshes: the CPU reference also appends its own far-shell
+        // repair geometry after surfaceVertexCount. Keep the original aggregate diagnostic
+        // above, and report the corresponding components independently here.
+        const std::size_t surfaceVertexCount=std::min<std::size_t>(
+            tile.expectedVertexCount,hwFirst.vertices.size());
+        const std::size_t warpSurfaceVertexCount=std::min<std::size_t>(
+            tile.expectedVertexCount,warp.vertices.size());
+        const std::vector<GpuVertex> gpuSurface(hwFirst.vertices.begin(),
+            hwFirst.vertices.begin()+surfaceVertexCount);
+        const std::vector<GpuVertex> gpuConnectorPanels(hwFirst.vertices.begin()+surfaceVertexCount,
+            hwFirst.vertices.end());
+        const std::vector<GpuVertex> warpSurface(warp.vertices.begin(),
+            warp.vertices.begin()+warpSurfaceVertexCount);
+        std::vector<GpuVertex> cpuCompleteMesh;cpuCompleteMesh.reserve(cpuMesh.vertices.size());
+        for(const MeshVertex& in:cpuMesh.vertices)
+        {
+            GpuVertex out{};
+            out.position[0]=static_cast<float>(in.position.x);
+            out.position[1]=static_cast<float>(in.position.y);
+            out.position[2]=static_cast<float>(in.position.z);
+            cpuCompleteMesh.push_back(out);
+        }
+        const DistanceStats baseToCpu=NearestDistance(gpuSurface,tile.cpuSurfaceVertices);
+        const DistanceStats cpuToBase=NearestDistance(tile.cpuSurfaceVertices,gpuSurface);
+        const double baseDirectedCount=static_cast<double>(baseToCpu.count+cpuToBase.count);
+        const double baseMean=baseDirectedCount
+            ?(baseToCpu.meanCm*baseToCpu.count+cpuToBase.meanCm*cpuToBase.count)/baseDirectedCount:0.0;
+        const float baseMaximum=std::max(baseToCpu.maxCm,cpuToBase.maxCm);
+        const DistanceStats panelsToCpuComplete=NearestDistance(gpuConnectorPanels,cpuCompleteMesh);
+        const DistanceStats gpuToCpuComplete=NearestDistance(hwFirst.vertices,cpuCompleteMesh);
+        const DistanceStats cpuCompleteToGpu=NearestDistance(cpuCompleteMesh,hwFirst.vertices);
+        const double completeDirectedCount=static_cast<double>(gpuToCpuComplete.count+cpuCompleteToGpu.count);
+        const double completeMean=completeDirectedCount
+            ?(gpuToCpuComplete.meanCm*gpuToCpuComplete.count
+                +cpuCompleteToGpu.meanCm*cpuCompleteToGpu.count)/completeDirectedCount:0.0;
+        const float completeMaximum=std::max(gpuToCpuComplete.maxCm,cpuCompleteToGpu.maxCm);
+        const DistanceStats gpuSurfaceToWarp=NearestDistance(gpuSurface,warpSurface);
+        const DistanceStats warpToGpuSurface=NearestDistance(warpSurface,gpuSurface);
+        const double surfaceWarpCount=static_cast<double>(gpuSurfaceToWarp.count+warpToGpuSurface.count);
+        const double surfaceWarpMean=surfaceWarpCount
+            ?(gpuSurfaceToWarp.meanCm*gpuSurfaceToWarp.count
+                +warpToGpuSurface.meanCm*warpToGpuSurface.count)/surfaceWarpCount:0.0;
+        const float surfaceWarpMaximum=std::max(gpuSurfaceToWarp.maxCm,warpToGpuSurface.maxCm);
+        const float baseVertexBoundCm=VertexQuantizationBoundCm(
+            gpuSurface,warpSurface,tile.cpuSurfaceVertices);
+        std::printf("[GpuVertexAgreementSplit] tile=(%d,%d,%d,lod%d) mc_vertices=%zu gpu_panel_vertices=%zu cpu_mesh_vertices=%zu cpu_surface_vertices=%zu mc_vs_cpu_surface_count=%llu mean_cm=%.8g max_cm=%.8g float_bound_cm=%.8g mc_vs_cpu_surface_bound_pass=%d panels_vs_cpu_complete_count=%llu mean_cm=%.8g max_cm=%.8g complete_meshes_count=%llu mean_cm=%.8g max_cm=%.8g mc_gpu_vs_warp_count=%llu mean_cm=%.8g max_cm=%.8g\n",
+            tile.key.coordinate.x,tile.key.coordinate.y,tile.key.coordinate.z,tile.key.lod,
+            surfaceVertexCount,gpuConnectorPanels.size(),cpuCompleteMesh.size(),tile.cpuSurfaceVertices.size(),
+            static_cast<unsigned long long>(baseToCpu.count+cpuToBase.count),baseMean,baseMaximum,
+            baseVertexBoundCm,baseMaximum<=baseVertexBoundCm,
+            static_cast<unsigned long long>(panelsToCpuComplete.count),
+            panelsToCpuComplete.meanCm,panelsToCpuComplete.maxCm,
+            static_cast<unsigned long long>(gpuToCpuComplete.count+cpuCompleteToGpu.count),
+            completeMean,completeMaximum,
+            static_cast<unsigned long long>(gpuSurfaceToWarp.count+warpToGpuSurface.count),
+            surfaceWarpMean,surfaceWarpMaximum);
         double fieldErrorSum=0.0,warpFieldErrorSum=0.0;
         float fieldErrorMax=0.0f,warpFieldErrorMax=0.0f;
         std::uint64_t fieldSignDiff=0,warpFieldSignDiff=0;
@@ -1893,11 +2029,16 @@ namespace
                 else ++fieldHaloSignDiff;
                 if(printedSignDiffSamples<8)
                 {
-                    std::printf("[FieldSignDifferenceSample] tile=(%d,%d,%d,lod%d) local=(%d,%d,%d) world=(%d,%d,%d) cpu=%.9g gpu=%.9g case_lattice=%d\n",
+                    const Vec3d worldPosition{
+                        double(tile.origin[0]+x*static_cast<int>(tile.step)),
+                        double(tile.origin[1]+y*static_cast<int>(tile.step)),
+                        double(tile.origin[2]+z*static_cast<int>(tile.step))};
+                    const float rawAnalytic=world.unclampedDensityAt(worldPosition);
+                    std::printf("[FieldSignDifferenceSample] tile=(%d,%d,%d,lod%d) local=(%d,%d,%d) world=(%d,%d,%d) cpu_packed=%.9g cpu_raw_analytic=%.9g gpu=%.9g case_lattice=%d\n",
                         tile.key.coordinate.x,tile.key.coordinate.y,tile.key.coordinate.z,tile.key.lod,
                         x,y,z,tile.origin[0]+x*static_cast<int>(tile.step),
                         tile.origin[1]+y*static_cast<int>(tile.step),
-                        tile.origin[2]+z*static_cast<int>(tile.step),tile.density[I],
+                        tile.origin[2]+z*static_cast<int>(tile.step),tile.density[I],rawAnalytic,
                         hwFirst.fieldDensity[I],x>=0&&x<=static_cast<int>(tile.cells)
                             &&y>=0&&y<=static_cast<int>(tile.cells)
                             &&z>=0&&z<=static_cast<int>(tile.cells));
@@ -1925,9 +2066,12 @@ namespace
         }
         const double fieldErrorMean=fieldCount?fieldErrorSum/fieldCount:0.0;
         const double warpFieldErrorMean=warpFieldCount?warpFieldErrorSum/warpFieldCount:0.0;
-        std::printf("[GpuTile] key=(%d,%d,%d) lod=%d cells=%u source_samples=%zu world_shapes=%zu world_bin_shape_ids=%zu protected_samples=%llu critical_lattice=%llu critical_cells=%llu cpu_vertices=%u gpu_vertices=%zu warp_vertices=%zu expected_gpu_vertices=%u gpu_indices_sequential=%d warp_indices_sequential=%d gpu_indexed_indirect_args=%d warp_indexed_indirect_args=%d d3d_buffer_mb=%.3f gpu_adapter_memory_mb=(%.2f,%.2f) budget_mb=%.2f warp_adapter_memory_info=%d gpu_field_ms=%.4f gpu_dispatch_ms=%.4f gpu_wall_ms=%.4f repeat_field_ms=%.4f repeat_dispatch_ms=%.4f warp_field_ms=%.4f warp_dispatch_ms=%.4f warp_wall_ms=%.4f\n",
+        const std::size_t nearAmbiguousSignOverrides=static_cast<std::size_t>(std::count_if(
+            farInput.nearPackedProxy.begin(),farInput.nearPackedProxy.end(),
+            [](std::uint32_t value){return (value&0x40000000u)!=0u;}));
+        std::printf("[GpuTile] key=(%d,%d,%d) lod=%d cells=%u source_samples=%zu near_ambiguous_sign_overrides=%zu world_shapes=%zu world_bin_shape_ids=%zu protected_samples=%llu critical_lattice=%llu critical_cells=%llu cpu_vertices=%u gpu_vertices=%zu warp_vertices=%zu expected_gpu_vertices=%u gpu_indices_sequential=%d warp_indices_sequential=%d gpu_indexed_indirect_args=%d warp_indexed_indirect_args=%d d3d_buffer_mb=%.3f gpu_adapter_memory_mb=(%.2f,%.2f) budget_mb=%.2f warp_adapter_memory_info=%d gpu_field_ms=%.4f gpu_dispatch_ms=%.4f gpu_wall_ms=%.4f repeat_field_ms=%.4f repeat_dispatch_ms=%.4f warp_field_ms=%.4f warp_dispatch_ms=%.4f warp_wall_ms=%.4f\n",
             tile.key.coordinate.x, tile.key.coordinate.y, tile.key.coordinate.z, tile.key.lod,
-            tile.cells, tile.density.size(),farInput.worldInput->shapes.size(),
+            tile.cells, tile.density.size(),nearAmbiguousSignOverrides,farInput.worldInput->shapes.size(),
             farInput.worldInput->spatialBinShapeIds.size(),
             static_cast<unsigned long long>(farInput.protectedSampleCount),static_cast<unsigned long long>(tile.criticalSamples),
             static_cast<unsigned long long>(tile.criticalCells), cpuMesh.groups.surfaceVertexCount,
@@ -2029,12 +2173,30 @@ namespace
         const std::filesystem::path dataPath = root / "data" / "owner_world.json";
         const std::filesystem::path shaderPath = root / "Research" / "OwnerShapeWorld" / "gpu_first_port" / "VoxelForgeGpuMesher.hlsl";
         bool wholeFar=false;
+        bool wholeNear=false;
         bool profileGpuKernels=false;
         int clipRadiusOverride=-1;
+        int tileLimit=0;
+        bool hasCameraOverride=false;
+        Vec3d cameraOverrideVoxels{};
         for(int i=1;i<argc;++i)
         {
             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])=="--profile-kernels")profileGpuKernels=true;
+            const std::string_view arg=argv[i];
+            if(arg.starts_with("--tile-limit="))
+                tileLimit=std::max(0,std::atoi(std::string(arg.substr(13)).c_str()));
+            if(arg.starts_with("--camera="))
+            {
+                float x=0.0f,y=0.0f,z=0.0f;
+                const std::string values(arg.substr(9));
+                if(std::sscanf(values.c_str(),"%f,%f,%f",&x,&y,&z)!=3)
+                    throw std::runtime_error("invalid --camera value: "+values);
+                cameraOverrideVoxels={static_cast<double>(x)*4.0,
+                    static_cast<double>(y)*4.0,static_cast<double>(z)*4.0};
+                hasCameraOverride=true;
+            }
         }
         for(int i=1;i<argc;++i)
         {
@@ -2058,7 +2220,10 @@ namespace
         if (!snapshot) throw std::runtime_error("snapshot load failed: " + error);
         World world(std::move(*snapshot));
         const WorldConfiguration config = world.configuration();
-        const Vec3d camera = config.viewerStartFound ? config.viewerStartVoxels : Vec3d{};
+        const Vec3d camera = hasCameraOverride?cameraOverrideVoxels:
+            (config.viewerStartFound ? config.viewerStartVoxels : Vec3d{});
+        std::printf("[GpuPortCamera] source=%s voxel=(%.3f,%.3f,%.3f)\n",
+            hasCameraOverride?"argument":"owner-start",camera.x,camera.y,camera.z);
         const std::vector<TileKey> desired = GetFullDesired(world, config, camera);
         PrintProcessMemory("after_snapshot");
         ShapeBuildRequest request;
@@ -2068,9 +2233,10 @@ namespace
         const auto compileStart = Clock::now();
         if (!world.compile(request, scheduler, error)) throw std::runtime_error("shape compile failed: " + error);
         const double compileMs = std::chrono::duration<double, std::milli>(Clock::now() - compileStart).count();
-        std::printf("[GpuPortWorld] desired_tiles=%zu camera_vox=(%.3f,%.3f,%.3f) full_resolution_lods=%d coarse_cells=%d shape_compile_ms=%.3f shape_stats_ms=(%.3f,%.3f,%.3f,%.3f) memory_limit_mb=8192\n",
+        std::printf("[GpuPortWorld] desired_tiles=%zu camera_vox=(%.3f,%.3f,%.3f) full_resolution_lods=%d coarse_cells=%d shape_bounds_min=(%d,%d,%d) shape_bounds_max=(%d,%d,%d) shape_compile_ms=%.3f shape_stats_ms=(%.3f,%.3f,%.3f,%.3f) memory_limit_mb=8192\n",
             desired.size(), camera.x, camera.y, camera.z, config.fullResolutionLodCount,
-            config.coarseTileCells, compileMs,
+            config.coarseTileCells,request.minVoxel.x,request.minVoxel.y,request.minVoxel.z,
+            request.maxVoxel.x,request.maxVoxel.y,request.maxVoxel.z,compileMs,
             world.buildStats().shapeSeconds * 1000.0, world.buildStats().intervalSeconds * 1000.0,
             world.buildStats().surfaceSeconds * 1000.0, world.buildStats().farShapeSeconds * 1000.0);
         PrintProcessMemory("after_shape_compile");
@@ -2095,7 +2261,7 @@ namespace
             static_cast<unsigned long long>(noiseRepeat.differingBits), noiseWarp.mean,
             noiseWarp.maximum, static_cast<unsigned long long>(noiseWarp.differingBits));
 
-        const std::vector<FarSelection> selected = SelectFarTiles(world, config, camera,wholeFar);
+        const std::vector<FarSelection> selected = SelectFarTiles(world, config, camera,wholeFar,wholeNear);
         if (selected.empty()) throw std::runtime_error("no far-field tiles were selected");
         bool allPass = noiseCpuGpu.differingBits == 0 && noiseRepeat.differingBits == 0
             && noiseWarp.differingBits == 0;
@@ -2110,14 +2276,16 @@ namespace
         std::uint64_t protectedThinSamples=0,thinFeatureCellsTotal=0,thinFeatureDiffGpuTotal=0,thinFeatureDiffWarpTotal=0;
         for (const FarSelection& selection : selected)
         {
+            if(tileLimit>0&&testedTiles>=static_cast<std::uint64_t>(tileLimit))break;
             const TileKey& key=selection.key;
             FarFieldTileInput farInputStorage;
             const FarFieldTileInput* farInput=&selection.input;
             double exportMs=selection.exportPreparationMs;
-            if(wholeFar)
+            if(wholeFar||wholeNear)
             {
                 const auto exportStart=Clock::now();
-                auto exported=world.exportFarFieldTile(key,error);
+                auto exported=wholeNear?world.exportGpuFieldTile(key,error)
+                    :world.exportFarFieldTile(key,error);
                 exportMs=std::chrono::duration<double,std::milli>(Clock::now()-exportStart).count();
                 if(!exported)
                 {
@@ -2126,6 +2294,10 @@ namespace
                     throw std::runtime_error("far field export failed: "+error);
                 }
                 farInputStorage=std::move(*exported);farInput=&farInputStorage;
+                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",
+                        wholeNear,farInput->nearDetailEnabled,farInput->cells,farInput->step);
             }
             TileGenerationOptions options;
             options.captureDensitySamples = true;
@@ -2158,7 +2330,7 @@ namespace
                     thinFeatureDiffGpuTotal+=gpuDiff;thinFeatureDiffWarpTotal+=warpDiff;
                 }
             }
-            const bool agreementPass = ReportTile(tile, *cpuMesh, gpu1, gpu2, warpResult,*farInput);
+            const bool agreementPass = ReportTile(world,tile, *cpuMesh, gpu1, gpu2, warpResult,*farInput);
             std::printf("[GpuFieldPipelineWall] key=(%d,%d,%d) lod=%d cpu_export_shape_sign_upload_prep_ms=%.4f cpu_generateTile_ms=%.4f cpu_reference_lattice_topology_check_ms=%.4f rx7900_field_gpu_ms=%.4f rx7900_mesh_plus_field_gpu_ms=%.4f rx7900_wall_upload_dispatch_readback_ms=%.4f warp_field_gpu_ms=%.4f warp_mesh_plus_field_wall_ms=%.4f\n",
                 key.coordinate.x, key.coordinate.y, key.coordinate.z, key.lod,
                 exportMs,cpuWallMs,inputWallMs,gpu1.fieldMs,gpu1.dispatchMs,gpu1.wallMs,
@@ -2198,8 +2370,8 @@ namespace
             protectedThinSamples+=farInput->protectedThinSampleCount;
             PrintProcessMemory("after_tile_pair");
         }
-        std::printf("[GpuFarPortSummary] tested_tiles=%llu mode=%s skipped_near_transition_tiles=%llu levels=2..8 cpu_export_shape_protected_prep_sum_ms=%.4f cpu_generateTile_sum_ms=%.4f cpu_reference_lattice_topology_check_sum_ms=%.4f rx7900_field_dispatch_sum_ms=%.4f rx7900_field_plus_meshing_sum_ms=%.4f warp_field_dispatch_sum_ms=%.4f topology_diff_cells_gpu=%llu topology_diff_cells_warp=%llu protected_exception_cells=%llu protected_exception_diff_gpu=%llu protected_exception_diff_warp=%llu thin_feature_cells=%llu thin_feature_diff_gpu=%llu thin_feature_diff_warp=%llu route_air_samples=%llu seal_override_samples=%llu thin_route_band_samples=%llu solid_overlay_samples=%llu whole_view_gpu_far=viewer_measurement_required viewer_direct_draw=integrated shell_survival=viewer_gpu_audit_required\n",
+        std::printf("[GpuFarPortSummary] tested_tiles=%llu mode=%s skipped_near_transition_tiles=%llu cpu_export_shape_protected_prep_sum_ms=%.4f cpu_generateTile_sum_ms=%.4f cpu_reference_lattice_topology_check_sum_ms=%.4f rx7900_field_dispatch_sum_ms=%.4f rx7900_field_plus_meshing_sum_ms=%.4f warp_field_dispatch_sum_ms=%.4f topology_diff_cells_gpu=%llu topology_diff_cells_warp=%llu protected_exception_cells=%llu protected_exception_diff_gpu=%llu protected_exception_diff_warp=%llu thin_feature_cells=%llu thin_feature_diff_gpu=%llu thin_feature_diff_warp=%llu route_air_samples=%llu seal_override_samples=%llu thin_route_band_samples=%llu solid_overlay_samples=%llu whole_view_gpu_far=viewer_measurement_required viewer_direct_draw=integrated shell_survival=viewer_gpu_audit_required\n",
-            static_cast<unsigned long long>(testedTiles),wholeFar?"whole-far-eligible":"sampled",
+            static_cast<unsigned long long>(testedTiles),wholeFar?"whole-far-eligible":(wholeNear?"whole-near-detail":"sampled"),
             static_cast<unsigned long long>(nearTransitionSkipped),
             cpuExportPrepSum,cpuTileWallSum,cpuFieldPrepSum,gpuFieldSum,
             gpuTotalSum,warpFieldSum,

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep6CaptureMatrix.ps1

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

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep6ColdInterleaved.ps1

@@ -0,0 +1,56 @@
+[CmdletBinding()]
+param([string]$ExecutableName='VoxelForgeGpuFieldStep6NearTimeline.exe')
+
+$ErrorActionPreference='Stop'
+$HostProcess=Get-Process -Id $PID
+$HostProcess.PriorityClass='BelowNormal'
+$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$ExecutablePath=Join-Path $StandaloneRoot ('bin\'+$ExecutableName)
+if([IO.Path]::GetFileName($ExecutableName)-ne$ExecutableName-or[IO.Path]::GetExtension($ExecutableName)-ne'.exe')
+{throw 'ExecutableName must be an executable filename.'}
+if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "Viewer not found: $ExecutablePath"}
+$Sequence=@(
+    @{Mode='cpu';Name='step6_final_cpu_A'},
+    @{Mode='gpu';Name='step6_final_gpu_A'},
+    @{Mode='gpu';Name='step6_final_gpu_B'},
+    @{Mode='cpu';Name='step6_final_cpu_B'},
+    @{Mode='cpu';Name='step6_final_cpu_C'},
+    @{Mode='gpu';Name='step6_final_gpu_C'}
+)
+foreach($Run in $Sequence)
+{
+    $Arguments=@('--clip-radius=3','--workers=8','--quit-after-clipmap',"--$($Run.Mode)-far")
+    $StdoutPath=Join-Path $PSScriptRoot ($Run.Name+'.stdout.log')
+    $StderrPath=Join-Path $PSScriptRoot ($Run.Name+'.stderr.log')
+    Write-Output "COLD_RUN mode=$($Run.Mode) command=`"$ExecutablePath`" $($Arguments -join ' ')"
+    for($Attempt=0;$Attempt-lt2;$Attempt++)
+    {
+        if($Attempt-gt0)
+        {Write-Host '0xC0000142 encountered; waiting 60 seconds before the single permitted retry.';Start-Sleep -Seconds 60}
+        Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
+        $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
+            -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
+        $PeakWorkingSet=0L;$PeakPrivateMemory=0L;$OverMemoryLimit=$false
+        while(-not$Child.HasExited)
+        {
+            $Child.Refresh();if($Child.HasExited){break}
+            $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+            $PeakWorkingSet=[Math]::Max($PeakWorkingSet,$Child.WorkingSet64)
+            $PeakPrivateMemory=[Math]::Max($PeakPrivateMemory,$Child.PrivateMemorySize64)
+            if($Child.PrivateMemorySize64-ge8GB)
+            {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+            Start-Sleep -Milliseconds 100
+        }
+        $Child.WaitForExit();$Child.Refresh()
+        if($Child.ExitCode-eq-1073741502-and$Attempt-eq0){continue}
+        if($Child.ExitCode-eq-1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
+        if($OverMemoryLimit){throw "Viewer stopped at the private-memory limit: $($Run.Name)"}
+        $LogText=Get-Content -LiteralPath $StdoutPath -Raw
+        if(-not$LogText.Contains('[InitialClipmapReady]'))
+        {throw "Clipmap was not ready in $($Run.Name); exit=$($Child.ExitCode)"}
+        Write-Output "RUN_EXIT mode=$($Run.Mode) log=$($Run.Name) pid=$($Child.Id) priority=BelowNormal affinity=0xFF00 peak_working_set_gb=$([math]::Round($PeakWorkingSet/1GB,2)) peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) exit=$($Child.ExitCode)"
+        if($Child.ExitCode-ne0){throw "Cold run returned $($Child.ExitCode): $($Run.Name)"}
+        break
+    }
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep6FarAgreement.ps1

@@ -0,0 +1,35 @@
+[CmdletBinding()]
+param()
+$ErrorActionPreference='Stop'
+$hostProcess=Get-Process -Id $PID
+$hostProcess.PriorityClass='BelowNormal'
+$hostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$standalone=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$exe=Join-Path $standalone 'bin\VoxelForgeGpuFieldStep6NearRunnerStep6.exe'
+$stdout=Join-Path $PSScriptRoot 'step6_whole_far_agreement.stdout.log'
+$stderr=Join-Path $PSScriptRoot 'step6_whole_far_agreement.stderr.log'
+$arguments=@('--whole-far','--clip-radius=3','--profile-kernels')
+for($attempt=0;$attempt -lt 2;$attempt++){
+    if($attempt -eq 1){Write-Output '0xC0000142 encountered; waiting 60 seconds before the single retry.';Start-Sleep -Seconds 60}
+    Remove-Item -LiteralPath $stdout,$stderr -Force -ErrorAction SilentlyContinue
+    $child=Start-Process -FilePath $exe -ArgumentList $arguments -WorkingDirectory $standalone `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $stdout -RedirectStandardError $stderr
+    $peakWs=0L;$peakPrivate=0L
+    while(-not $child.HasExited){
+        $child.Refresh()
+        if($child.HasExited){break}
+        $child.PriorityClass='BelowNormal'
+        $child.ProcessorAffinity=[IntPtr]0xFF00
+        $peakWs=[Math]::Max($peakWs,$child.WorkingSet64)
+        $peakPrivate=[Math]::Max($peakPrivate,$child.PrivateMemorySize64)
+        if($child.PrivateMemorySize64 -ge 8GB){$child.Kill();$child.WaitForExit();throw 'far agreement exceeded 8 GB private memory'}
+        Start-Sleep -Milliseconds 100
+    }
+    $child.WaitForExit()
+    if($child.ExitCode -eq -1073741502 -and $attempt -eq 0){continue}
+    if($child.ExitCode -eq -1073741502){throw '0xC0000142 repeated after its single retry'}
+    $log=Get-Content -LiteralPath $stdout -Raw
+    if(-not $log.Contains('[GpuFarPortSummary]') -or -not $log.Contains('[GpuPortStatus]')){throw "far agreement did not finish; exit=$($child.ExitCode)"}
+    Write-Output "RUN_EXIT pid=$($child.Id) priority=BelowNormal affinity=0xFF00 peak_working_set_gb=$([Math]::Round($peakWs/1GB,2)) peak_private_gb=$([Math]::Round($peakPrivate/1GB,2)) stdout=$stdout stderr=$stderr"
+    exit 0
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep6FarPacked.ps1

@@ -0,0 +1,35 @@
+[CmdletBinding()]
+param()
+$ErrorActionPreference='Stop'
+$hostProcess=Get-Process -Id $PID
+$hostProcess.PriorityClass='BelowNormal'
+$hostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$standalone=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$exe=Join-Path $standalone 'bin\VoxelForgeGpuFieldStep6NearPackedRunner.exe'
+$stdout=Join-Path $PSScriptRoot 'step6_whole_far_packed.stdout.log'
+$stderr=Join-Path $PSScriptRoot 'step6_whole_far_packed.stderr.log'
+$arguments=@('--whole-far','--clip-radius=3','--profile-kernels')
+for($attempt=0;$attempt -lt 2;$attempt++){
+    if($attempt -eq 1){Write-Output '0xC0000142 encountered; waiting 60 seconds before one retry.';Start-Sleep -Seconds 60}
+    Remove-Item -LiteralPath $stdout,$stderr -Force -ErrorAction SilentlyContinue
+    $child=Start-Process -FilePath $exe -ArgumentList $arguments -WorkingDirectory $standalone `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $stdout -RedirectStandardError $stderr
+    $peakWs=0L;$peakPrivate=0L
+    while(-not $child.HasExited){
+        $child.Refresh()
+        if($child.HasExited){break}
+        $child.PriorityClass='BelowNormal'
+        $child.ProcessorAffinity=[IntPtr]0xFF00
+        $peakWs=[Math]::Max($peakWs,$child.WorkingSet64)
+        $peakPrivate=[Math]::Max($peakPrivate,$child.PrivateMemorySize64)
+        if($child.PrivateMemorySize64 -ge 8GB){$child.Kill();$child.WaitForExit();throw 'far audit exceeded 8 GB private memory'}
+        Start-Sleep -Milliseconds 100
+    }
+    $child.WaitForExit()
+    if($child.ExitCode -eq -1073741502 -and $attempt -eq 0){continue}
+    if($child.ExitCode -eq -1073741502){throw '0xC0000142 repeated after one retry'}
+    $log=Get-Content -LiteralPath $stdout -Raw
+    if(-not $log.Contains('[GpuFarPortSummary]') -or -not $log.Contains('[GpuPortStatus]')){throw "far audit did not finish; exit=$($child.ExitCode)"}
+    Write-Output "RUN_EXIT pid=$($child.Id) priority=BelowNormal affinity=0xFF00 peak_working_set_gb=$([Math]::Round($peakWs/1GB,2)) peak_private_gb=$([Math]::Round($peakPrivate/1GB,2)) stdout=$stdout stderr=$stderr"
+    exit 0
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep6GpuCrossing.ps1

@@ -0,0 +1,54 @@
+[CmdletBinding()]
+param(
+    [string]$ExecutableName='VoxelForgeGpuFieldStep6NearTimeline.exe',
+    [string]$LogStem='step6_gpu_crossing',
+    [int]$FlightSpeedMps=16
+)
+
+$ErrorActionPreference='Stop'
+if([IO.Path]::GetFileName($ExecutableName)-ne$ExecutableName-or[IO.Path]::GetExtension($ExecutableName)-ne'.exe')
+{throw 'ExecutableName must be an executable filename.'}
+if([IO.Path]::GetFileName($LogStem)-ne$LogStem){throw 'LogStem must be a filename stem.'}
+$HostProcess=Get-Process -Id $PID
+$HostProcess.PriorityClass='BelowNormal'
+$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$ExecutablePath=Join-Path $StandaloneRoot ('bin\'+$ExecutableName)
+$StdoutPath=Join-Path $PSScriptRoot ($LogStem+'.stdout.log')
+$StderrPath=Join-Path $PSScriptRoot ($LogStem+'.stderr.log')
+if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "Viewer not found: $ExecutablePath"}
+$Arguments=@('--gpu-far','--clip-radius=3','--workers=8','--test-one-chunk-crossing',"--flight-speed-mps=$FlightSpeedMps")
+Write-Output "VIEWER_COMMAND `"$ExecutablePath`" $($Arguments -join ' ')"
+for($Attempt=0;$Attempt -lt 2;$Attempt++)
+{
+    if($Attempt-gt 0)
+    {
+        Write-Host '0xC0000142 encountered; waiting 60 seconds before the single permitted retry.'
+        Start-Sleep -Seconds 60
+    }
+    Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
+    $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
+    $PeakWorkingSet=0L;$PeakPrivateMemory=0L;$OverMemoryLimit=$false
+    while(-not$Child.HasExited)
+    {
+        $Child.Refresh()
+        if($Child.HasExited){break}
+        $Child.PriorityClass='BelowNormal'
+        $Child.ProcessorAffinity=[IntPtr]0xFF00
+        $PeakWorkingSet=[Math]::Max($PeakWorkingSet,$Child.WorkingSet64)
+        $PeakPrivateMemory=[Math]::Max($PeakPrivateMemory,$Child.PrivateMemorySize64)
+        if($Child.PrivateMemorySize64-ge8GB)
+        {$OverMemoryLimit=$true;$Child.Kill();$Child.WaitForExit();break}
+        Start-Sleep -Milliseconds 100
+    }
+    $Child.WaitForExit();$Child.Refresh()
+    if($Child.ExitCode-eq-1073741502-and$Attempt-eq0){continue}
+    if($Child.ExitCode-eq-1073741502){throw '0xC0000142 repeated after its single retry; stopping.'}
+    if($OverMemoryLimit){throw "Viewer stopped at the private-memory limit; stdout=$StdoutPath stderr=$StderrPath"}
+    $LogText=Get-Content -LiteralPath $StdoutPath -Raw
+    $Complete=$LogText.Contains('[FlightBenchmarkComplete]')
+    Write-Output "RUN_EXIT pid=$($Child.Id) priority=BelowNormal affinity=0xFF00 peak_working_set_gb=$([math]::Round($PeakWorkingSet/1GB,2)) peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) exit_code=$($Child.ExitCode) crossing_complete=$Complete stdout=$StdoutPath stderr=$StderrPath"
+    if(-not$Complete){exit 4}
+    exit $Child.ExitCode
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep6NearAgreement.ps1

@@ -0,0 +1,53 @@
+[CmdletBinding()]
+param()
+
+$ErrorActionPreference = 'Stop'
+$HostProcess = Get-Process -Id $PID
+$HostProcess.PriorityClass = 'BelowNormal'
+$HostProcess.ProcessorAffinity = [IntPtr]0xFF00
+$Standalone = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$Executable = Join-Path $Standalone 'bin\VoxelForgeGpuFieldStep6NearRunnerStep6.exe'
+$Stdout = Join-Path $PSScriptRoot 'step6_whole_near_agreement_fixed.stdout.log'
+$Stderr = Join-Path $PSScriptRoot 'step6_whole_near_agreement_fixed.stderr.log'
+$Arguments = @('--whole-near','--clip-radius=3','--profile-kernels')
+if (-not (Test-Path -LiteralPath $Executable -PathType Leaf)) { throw "Executable not found: $Executable" }
+Write-Output "STEP6_GPU_NEAR_COMMAND: $Executable $($Arguments -join ' ')"
+
+for ($Attempt = 0; $Attempt -lt 2; ++$Attempt) {
+    if ($Attempt -eq 1) {
+        Write-Output '0xC0000142 encountered; waiting 60 seconds before the permitted retry.'
+        Start-Sleep -Seconds 60
+    }
+    Remove-Item -LiteralPath $Stdout,$Stderr -Force -ErrorAction SilentlyContinue
+    $Child = Start-Process -FilePath $Executable -ArgumentList $Arguments -WorkingDirectory $Standalone `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $Stdout -RedirectStandardError $Stderr
+    $PeakWorkingSet = 0L
+    $PeakPrivate = 0L
+    $OverLimit = $false
+    while (-not $Child.HasExited) {
+        $Child.Refresh()
+        if ($Child.HasExited) { break }
+        $Child.PriorityClass = 'BelowNormal'
+        $Child.ProcessorAffinity = [IntPtr]0xFF00
+        $PeakWorkingSet = [Math]::Max($PeakWorkingSet,$Child.WorkingSet64)
+        $PeakPrivate = [Math]::Max($PeakPrivate,$Child.PrivateMemorySize64)
+        if ($Child.PrivateMemorySize64 -ge 8GB) {
+            $OverLimit = $true
+            $Child.Kill()
+            $Child.WaitForExit()
+            break
+        }
+        Start-Sleep -Milliseconds 100
+    }
+    $Child.WaitForExit()
+    $Child.Refresh()
+    if ($Child.ExitCode -eq -1073741502 -and $Attempt -eq 0) { continue }
+    if ($Child.ExitCode -eq -1073741502) { throw '0xC0000142 repeated after its single retry; stopping.' }
+    if ($OverLimit) { throw "Near agreement stopped at the 8 GB process memory limit; stdout=$Stdout stderr=$Stderr" }
+    $Log = Get-Content -LiteralPath $Stdout -Raw
+    if (-not $Log.Contains('[GpuFarPortSummary]') -or -not $Log.Contains('[GpuPortStatus]')) {
+        throw "Near agreement did not finish; exit=$($Child.ExitCode) stdout=$Stdout stderr=$Stderr"
+    }
+    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($PeakPrivate/1GB,2)) stdout=$Stdout stderr=$Stderr"
+    exit 0
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep6NearDiag.ps1

@@ -0,0 +1,30 @@
+[CmdletBinding()]
+param()
+$ErrorActionPreference='Stop'
+$hostProcess=Get-Process -Id $PID
+$hostProcess.PriorityClass='BelowNormal'
+$hostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$standalone=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$exe=Join-Path $standalone 'bin\VoxelForgeGpuFieldStep6NearDiag.exe'
+$stdout=Join-Path $PSScriptRoot 'step6_near_diag.stdout.log'
+$stderr=Join-Path $PSScriptRoot 'step6_near_diag.stderr.log'
+$arguments=@('--whole-near','--clip-radius=3','--tile-limit=1','--profile-kernels')
+for($attempt=0;$attempt -lt 2;$attempt++){
+    if($attempt -eq 1){Start-Sleep -Seconds 60}
+    Remove-Item -LiteralPath $stdout,$stderr -Force -ErrorAction SilentlyContinue
+    $child=Start-Process -FilePath $exe -ArgumentList $arguments -WorkingDirectory $standalone `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $stdout -RedirectStandardError $stderr
+    while(-not $child.HasExited){
+        $child.Refresh()
+        if($child.HasExited){break}
+        $child.PriorityClass='BelowNormal'
+        $child.ProcessorAffinity=[IntPtr]0xFF00
+        if($child.PrivateMemorySize64 -ge 8GB){$child.Kill();$child.WaitForExit();throw 'diagnostic exceeded 8 GB private memory'}
+        Start-Sleep -Milliseconds 100
+    }
+    $child.WaitForExit()
+    if($child.ExitCode -eq -1073741502 -and $attempt -eq 0){continue}
+    if($child.ExitCode -eq -1073741502){throw '0xC0000142 repeated after its single retry'}
+    Write-Output "RUN_EXIT pid=$($child.Id) exit=$($child.ExitCode) stdout=$stdout stderr=$stderr"
+    exit $child.ExitCode
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep6NearDiagFull.ps1

@@ -0,0 +1,33 @@
+[CmdletBinding()]
+param()
+$ErrorActionPreference='Stop'
+$hostProcess=Get-Process -Id $PID
+$hostProcess.PriorityClass='BelowNormal'
+$hostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$standalone=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$exe=Join-Path $standalone 'bin\VoxelForgeGpuFieldStep6NearDiag2.exe'
+$stdout=Join-Path $PSScriptRoot 'step6_near_diag_full.stdout.log'
+$stderr=Join-Path $PSScriptRoot 'step6_near_diag_full.stderr.log'
+$arguments=@('--whole-near','--clip-radius=3','--tile-limit=960','--profile-kernels')
+for($attempt=0;$attempt -lt 2;$attempt++){
+    if($attempt -eq 1){Write-Output '0xC0000142 encountered; waiting 60 seconds before one retry.';Start-Sleep -Seconds 60}
+    Remove-Item -LiteralPath $stdout,$stderr -Force -ErrorAction SilentlyContinue
+    $child=Start-Process -FilePath $exe -ArgumentList $arguments -WorkingDirectory $standalone `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $stdout -RedirectStandardError $stderr
+    $peakWs=0L;$peakPrivate=0L
+    while(-not $child.HasExited){
+        $child.Refresh()
+        if($child.HasExited){break}
+        $child.PriorityClass='BelowNormal'
+        $child.ProcessorAffinity=[IntPtr]0xFF00
+        $peakWs=[Math]::Max($peakWs,$child.WorkingSet64)
+        $peakPrivate=[Math]::Max($peakPrivate,$child.PrivateMemorySize64)
+        if($child.PrivateMemorySize64 -ge 8GB){$child.Kill();$child.WaitForExit();throw 'near diagnostic exceeded 8 GB private memory'}
+        Start-Sleep -Milliseconds 100
+    }
+    $child.WaitForExit()
+    if($child.ExitCode -eq -1073741502 -and $attempt -eq 0){continue}
+    if($child.ExitCode -eq -1073741502){throw '0xC0000142 repeated after one retry'}
+    Write-Output "RUN_EXIT pid=$($child.Id) priority=BelowNormal affinity=0xFF00 peak_working_set_gb=$([Math]::Round($peakWs/1GB,2)) peak_private_gb=$([Math]::Round($peakPrivate/1GB,2)) stdout=$stdout stderr=$stderr"
+    exit $child.ExitCode
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep6NearMismatchProbe.ps1

@@ -0,0 +1,53 @@
+[CmdletBinding()]
+param()
+
+$ErrorActionPreference = 'Stop'
+$HostProcess = Get-Process -Id $PID
+$HostProcess.PriorityClass = 'BelowNormal'
+$HostProcess.ProcessorAffinity = [IntPtr]0xFF00
+$Standalone = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$Executable = Join-Path $Standalone 'bin\VoxelForgeGpuFieldStep6NearProbe8.exe'
+$Stdout = Join-Path $PSScriptRoot 'step6_near_mismatch_probe.stdout.log'
+$Stderr = Join-Path $PSScriptRoot 'step6_near_mismatch_probe.stderr.log'
+$Arguments = @('--whole-near','--clip-radius=3','--tile-limit=3','--profile-kernels')
+if (-not (Test-Path -LiteralPath $Executable -PathType Leaf)) { throw "Executable not found: $Executable" }
+Write-Output "STEP6_GPU_NEAR_PROBE_COMMAND: $Executable $($Arguments -join ' ')"
+
+for ($Attempt = 0; $Attempt -lt 2; ++$Attempt) {
+    if ($Attempt -eq 1) {
+        Write-Output '0xC0000142 encountered; waiting 60 seconds before the permitted retry.'
+        Start-Sleep -Seconds 60
+    }
+    Remove-Item -LiteralPath $Stdout,$Stderr -Force -ErrorAction SilentlyContinue
+    $Child = Start-Process -FilePath $Executable -ArgumentList $Arguments -WorkingDirectory $Standalone `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $Stdout -RedirectStandardError $Stderr
+    $PeakWorkingSet = 0L
+    $PeakPrivate = 0L
+    $OverLimit = $false
+    while (-not $Child.HasExited) {
+        $Child.Refresh()
+        if ($Child.HasExited) { break }
+        $Child.PriorityClass = 'BelowNormal'
+        $Child.ProcessorAffinity = [IntPtr]0xFF00
+        $PeakWorkingSet = [Math]::Max($PeakWorkingSet,$Child.WorkingSet64)
+        $PeakPrivate = [Math]::Max($PeakPrivate,$Child.PrivateMemorySize64)
+        if ($Child.PrivateMemorySize64 -ge 8GB) {
+            $OverLimit = $true
+            $Child.Kill()
+            $Child.WaitForExit()
+            break
+        }
+        Start-Sleep -Milliseconds 100
+    }
+    $Child.WaitForExit()
+    $Child.Refresh()
+    if ($Child.ExitCode -eq -1073741502 -and $Attempt -eq 0) { continue }
+    if ($Child.ExitCode -eq -1073741502) { throw '0xC0000142 repeated after its single retry; stopping.' }
+    if ($OverLimit) { throw "Near mismatch probe stopped at the 8 GB process memory limit; stdout=$Stdout stderr=$Stderr" }
+    $Log = Get-Content -LiteralPath $Stdout -Raw
+    if (-not $Log.Contains('[GpuFarPortSummary]') -or -not $Log.Contains('[GpuPortStatus]')) {
+        throw "Near mismatch probe did not finish; exit=$($Child.ExitCode) stdout=$Stdout stderr=$Stderr"
+    }
+    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($PeakPrivate/1GB,2)) stdout=$Stdout stderr=$Stderr"
+    exit 0
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep6NearPacked.ps1

@@ -0,0 +1,37 @@
+[CmdletBinding()]
+param(
+    [string]$Camera='',
+    [string]$LogStem='step6_whole_near_packed'
+)
+$ErrorActionPreference='Stop'
+$hostProcess=Get-Process -Id $PID
+$hostProcess.PriorityClass='BelowNormal'
+$hostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$standalone=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$exe=Join-Path $standalone 'bin\VoxelForgeGpuFieldStep6NearPackedRunner.exe'
+$stdout=Join-Path $PSScriptRoot ($LogStem+'.stdout.log')
+$stderr=Join-Path $PSScriptRoot ($LogStem+'.stderr.log')
+$arguments=@('--whole-near','--clip-radius=3','--tile-limit=960','--profile-kernels')
+if($Camera){$arguments+=("--camera=$Camera")}
+for($attempt=0;$attempt -lt 2;$attempt++){
+    if($attempt -eq 1){Write-Output '0xC0000142 encountered; waiting 60 seconds before one retry.';Start-Sleep -Seconds 60}
+    Remove-Item -LiteralPath $stdout,$stderr -Force -ErrorAction SilentlyContinue
+    $child=Start-Process -FilePath $exe -ArgumentList $arguments -WorkingDirectory $standalone `
+        -WindowStyle Hidden -PassThru -RedirectStandardOutput $stdout -RedirectStandardError $stderr
+    $peakWs=0L;$peakPrivate=0L
+    while(-not $child.HasExited){
+        $child.Refresh()
+        if($child.HasExited){break}
+        $child.PriorityClass='BelowNormal'
+        $child.ProcessorAffinity=[IntPtr]0xFF00
+        $peakWs=[Math]::Max($peakWs,$child.WorkingSet64)
+        $peakPrivate=[Math]::Max($peakPrivate,$child.PrivateMemorySize64)
+        if($child.PrivateMemorySize64 -ge 8GB){$child.Kill();$child.WaitForExit();throw 'near audit exceeded 8 GB private memory'}
+        Start-Sleep -Milliseconds 100
+    }
+    $child.WaitForExit()
+    if($child.ExitCode -eq -1073741502 -and $attempt -eq 0){continue}
+    if($child.ExitCode -eq -1073741502){throw '0xC0000142 repeated after one retry'}
+    Write-Output "RUN_EXIT pid=$($child.Id) priority=BelowNormal affinity=0xFF00 peak_working_set_gb=$([Math]::Round($peakWs/1GB,2)) peak_private_gb=$([Math]::Round($peakPrivate/1GB,2)) stdout=$stdout stderr=$stderr"
+    exit $child.ExitCode
+}

Standalone/Research/OwnerShapeWorld/gpu_field_port/STEP6_REPORT.md

@@ -0,0 +1,117 @@
+# VoxelForgeCore GPU port — step 6
+
+## Result
+
+The viewer now evaluates and meshes all 960 LOD0/1 tiles on the GPU, including the near detail band, relief, root refinement, and normals. It draws those output buffers directly; collision and CPU field queries remain on the CPU. Across three interleaved cold runs the GPU hybrid took **4.504 s**, against **2.506 s** for all CPU, so this port does not yet improve whole-view cold readiness. The GPU kernels are no longer the critical path.
+
+The remaining CPU/GPU vertex outlier is confined to far connector repair panels. Ordinary marching-cubes geometry and the complete near audit meet the millimetre target. The far-panel exception is quantified below; it is not used for collision and it does not change protected/thin topology.
+
+## Field split and shared code
+
+`VoxelForgeFieldShared.h` supplies the C++ and HLSL versions of the field math used here: primitive distances, profiled floors, blends, integer noise hashing, relief, and the shared root-refinement rule. The 4,096 shared-math probes matched CPU, RX 7900 XTX, and WARP bit-for-bit. The GPU evaluates ordinary rock/air from the analytic field. CPU uploads the world shape records and spatial bins once (18,401 shapes, 422,754 bins, 679,503 bin-shape IDs; 30.523 MiB), then per-tile keys and sparse route/seal/thin-feature exceptions. It does not upload a full CPU sign volume. A sparse one-quantum near-surface override handles the CPU quantizer's sign at 2,688 samples of 41,160,000 near lattice samples (0.00653%).
+
+The CPU still owns layout compilation, bin construction, exception discovery, and collision queries. These are not in the shared header. In the viewer, LOD0 edge midpoints are evaluated from the analytic HLSL field when the linear root is in the CPU refinement interval; the GPU also writes an unclamped near-density scratch field for normals where the meshing density is clamped. Audit runs separately compare against the CPU reference. GPU mesh output is render-only.
+
+## Cold timeline and critical path
+
+Command for the three paired runs:
+
+```text
+VoxelForgeGpuFieldStep6NearTimeline.exe --clip-radius=3 --workers=8 --quit-after-clipmap --cpu-far
+VoxelForgeGpuFieldStep6NearTimeline.exe --clip-radius=3 --workers=8 --quit-after-clipmap --gpu-far
+```
+
+The wrapper was `RunGpuFarViewer.ps1`, with CPU and GPU modes interleaved A/B/C. All launches used BelowNormal priority, affinity `0xFF00`, and eight workers. The viewer waits for completion queries without waiting inside each tile submission loop.
+
+| Cold stage | Measured result | What it waits on |
+|---|---:|---|
+| Owner data load | 13.8 ms | Input files and owner data |
+| D3D viewer/device initialization | 68.5 ms | Device and swap-chain setup |
+| Configure/prepare gap before shape compile | about 1.13 s | Owner resolution/parsing, parity checks, and `PrepareStandaloneDensityWorld`; this interval is not separately instrumented |
+| Shape compile | 590 ms wall; detail-line portion 368 ms wall / 2,891 ms worker sum | Compiled shape and detail-band data |
+| Persistent output pool | 64.5 ms average | 2,074 reserved GPU tile slots |
+| GPU input preparation | 228 ms wall / 1.756 s worker sum | Per-tile sparse inputs and GPU batch setup; world shape/bin data uploaded once |
+| GPU upload/enqueue | 94 ms summed across tiles | 66 batches, up to 32 tiles each |
+| GPU dispatch timestamps | 140.1 ms total / 0.06755 ms per tile | 2,074 tiles; maximum batch 9.9 ms |
+| First long `Present(1)` | about 2.10 s in the representative trace | Main thread blocks inside `Present(sync=1)`; instrumentation identifies the call, not the internal D3D cause |
+| Near detail ready | 4.358 s average | All 960 LOD0/1 GPU or CPU tiles |
+| Whole clipmap ready | 4.504 s average | All 2,074 GPU tiles, 152 CPU fallback tiles, and uploads |
+
+On the representative GPU trace, the queue envelope was 2.381 s, with 140 ms GPU-busy and an estimated 2.241 s idle-or-poll. Shape-detail workers averaged about 7.9 of 8 cores during their 368 ms wall interval; input export averaged about 7.7 of 8 cores. The serial wait is therefore not field-kernel execution: cold readiness waits on CPU world setup and shape compilation, GPU input preparation, the long `Present` call, then outstanding asynchronous tile completions and the CPU fallback uploads. The internal cause of the `Present` block remains unresolved. Near readiness waits for the 960 LOD0/1 tiles; whole readiness waits for the complete resident clipmap.
+
+| Interleaved cold run | All CPU | GPU hybrid |
+|---|---:|---:|
+| A | 2.552 s | 4.553 s |
+| B | 2.482 s | 4.548 s |
+| C | 2.484 s | 4.411 s |
+| Mean | **2.506 s** | **4.504 s** |
+
+Hybrid is 1.998 s (79.7%) slower in this cold test. GPU field/mesh work is fast; CPU export and startup/present costs erase the saved CPU tile work. The reported cold GPU time is D3D timestamp work only, while CPU and wall-clock numbers include their full measured stages.
+
+### Kernel profile
+
+The standalone full-near passage audit measured 960 LOD0/1 tiles:
+
+| Kernel | RX 7900 XTX GPU time |
+|---|---:|
+| Field | 27.50 ms |
+| Classify | 5.00 ms |
+| Scan | 8.43 ms |
+| Emit | 11.71 ms |
+| Connector rings | 48.68 ms |
+| Connector panels | 0 ms |
+| Total field plus mesh | 101.57 ms |
+
+It evaluated 24.49 million shape distances over 41.16 million lattice samples (mean 0.595 shape distances/sample, maximum 8). A second near camera measured 0.667/sample, maximum 11. The separate 1,114-tile far audit measured 113.60 ms total: field 15.61 ms, classify 7.01 ms, scan 4.22 ms, emit 26.80 ms, connector rings 58.07 ms, connector panels 1.61 ms. Far field shape-distance count averaged 0.1168/sample (maximum 9). These audit totals are per-tile timestamp sums, not whole-view cold wall time.
+
+## Agreement and vertex outliers
+
+| Region | Tiles | CPU/GPU sign / topology differences | WARP sign / topology differences | Protected/thin audit |
+|---|---:|---:|---:|---|
+| Near, ordinary camera | 960 | 0 / 0 | 0 / 0 | 229,376 protected cells; 0 differences |
+| Near, passage camera | 960 | 0 / 0 | 0 / 0 | 167,858 protected cells and 35,612 thin-feature cells; 0 differences on either implementation |
+| Far eligible | 1,114 | 0 / 0 | 0 / 0 | 127,401 protected cells and 353 thin-feature cells; 0 differences on either implementation |
+
+GPU repeats were bit-identical in the near and far audits. The HLSL shared-field probes were bit-identical across CPU, GPU, and WARP. WARP here is Microsoft's software D3D implementation. The GPU decides ordinary signs; route air, seals, and thin route features are supplied as sparse CPU exceptions.
+
+Near vertex agreement in the passage audit, using the audit's symmetric nearest-vertex metric:
+
+- CPU/GPU: mean **0.0000067 mm**, maximum **0.0049 mm** across 2,138,731 directed comparisons.
+- GPU/WARP: mean **0.0000026 mm**, maximum **0.0049 mm** across 3,633,900 directed comparisons.
+- Ordinary far marching-cubes vertices: maximum CPU/GPU distance **0.3125 mm**; all tested ordinary MC components pass their quantization bound.
+
+Across the 1,114-tile far audit, the aggregate has a **64.5602 cm** maximum nearest-vertex distance (mean 1.1524 mm; 247 of 901 nonempty tiles exceed their LOD quantization bound). It is not a large-world float precision or edge-root issue. The component split places the maximum entirely in GPU connector repair-panel vertices; ordinary marching-cubes vertices in that same tile differ by at most 0.3125 mm. The maximum is at far LOD8 tile `(1,2,-1)`, whose cells are 4 m across and whose tile spans 64 m. GPU emits connector repair panels over the intersecting connector region; the CPU reference emits sparse repair panels only where its marching-cubes shell misses the connector. The two panel sets therefore have different vertices and tessellation. This diagnostic is nearest-vertex distance, not a surface Hausdorff distance.
+
+This is a bounded exception only for the coarsest far visual connector repair: the measured maximum is under one LOD8 cell, protected/thin topology is identical, the far-survival audit passes, and the panels are never used for collision or near rendering. It does **not** establish millimetre agreement for the far connector panels. If those panels are used closer than the far LOD or require a strict geometric bound, the CPU's sparse repair-region plan still needs to be shared with or reproduced by the GPU. GPU/WARP far mesh vertices have mean 0.0047 mm and maximum 8.22 mm under the same nearest-vertex diagnostic; topology remains identical.
+
+## Survival, crossing, and frame slicing
+
+The final GPU-mesh far-survival audit passed: 7,200/7,200 boundary surfaces, 10,976/10,976 connectors, 178,351 connector samples, 710,047 connector rays, zero missing connector samples/rays, and zero connector holes. It recorded 50,019 extra ray intersections beyond each ray's first hit; no bad-shell triangles were removed. `GpuFarSurvivalAudit` confirms the source was GPU meshes with field readback enabled.
+
+The updated one-crossing run used:
+
+```text
+VoxelForgeGpuFieldStep6NearTimeline.exe --gpu-far --clip-radius=3 --workers=8 --test-one-chunk-crossing --flight-speed-mps=16
+```
+
+At 16 m/s, the nearest LOD0 tile was ready **22.931 ms** after the crossing; crossing work finished in 24 ms wall time with 7 ms summed worker CPU time. This is one crossing sample, so p50/p95/max all equal 22.931 ms. The run completed with all 2,226 resident tiles present and uploads drained.
+
+Using the measured 0.06755 ms/tile dispatch mean, a sliced GPU workload fits about 7 tiles in a 0.5 ms GPU slice (296 frames / 3.29 s for 2,074 tiles), or 14 tiles in a 1.0 ms slice (148 frames / 1.64 s). A current 32-tile batch can take about 9.9 ms, close to the 11.111 ms 90 Hz frame budget; batches should be split for VR pacing. These figures cover GPU dispatch only, not CPU preparation or submission.
+
+## GTX 1060-class estimate
+
+The [RX 7900 XTX product specifications](https://www.amd.com/en/products/graphics/desktops/radeon/7000-series/amd-radeon-rx-7900xtx.html) list 61.4 TFLOPS FP32; at the user's approximate 4.4 TFLOPS for a GTX 1060-class card, the simple throughput ratio is about 14:1. Scaling the measured 140 ms dispatch total gives a 1.95 s idealized estimate. Pascal's slower 32-bit integer multiply and driver effects make **about 2–3 s of GPU compute** for these 2,074 tiles a more plausible estimate, not a measurement. With the same measured CPU setup/export and first-use D3D costs, estimate roughly **6–7 s cold hybrid**. This is not an estimate for an optimized GTX implementation.
+
+## Capture/look audit
+
+All six fixed CPU/GPU pairs and four remote-yaw pairs were refreshed after the near normal fix. The full pixel comparison is in [step6_capture_comparison.md](step6_capture_comparison.md). Fixed views had RGB MAE 0.61–1.61; far fixed view MAE was 0.79. Passage had 8.95% of pixels above 8 channel values. Remote views differed more: MAE 3.37–6.84, with 5.27–9.85% of pixels above 32 channel values. These are image differences, not topology failures. The remote/far image differences are compatible with the still-different connector repair coverage, though the captures do not isolate it as the sole cause.
+
+Fixed pairs: [chamber CPU](step6_cpu_chamber.png) / [GPU](step6_gpu_chamber.png), [tunnel CPU](step6_cpu_tunnel.png) / [GPU](step6_gpu_tunnel.png), [passage CPU](step6_cpu_passage.png) / [GPU](step6_gpu_passage.png), [far CPU](step6_cpu_far.png) / [GPU](step6_gpu_far.png), [floor close-up CPU](step6_cpu_passage_floor_closeup.png) / [GPU](step6_gpu_passage_floor_closeup.png), [tunnel bend CPU](step6_cpu_tunnel_bend.png) / [GPU](step6_gpu_tunnel_bend.png).
+
+Remote pairs: [yaw 0 CPU](step6_cpu_remote_yaw0.png) / [GPU](step6_gpu_remote_yaw0.png), [yaw 90 CPU](step6_cpu_remote_yaw90.png) / [GPU](step6_gpu_remote_yaw90.png), [yaw 180 CPU](step6_cpu_remote_yaw180.png) / [GPU](step6_gpu_remote_yaw180.png), [yaw 270 CPU](step6_cpu_remote_yaw270.png) / [GPU](step6_gpu_remote_yaw270.png).
+
+## Reference and process constraints
+
+`CompareFarCpuHashes.ps1` passed against the saved `566d1f3` reference: **44 hash fields, 0 differences**. Runtime workers were 8; viewer and audit launchers set BelowNormal priority and affinity `0xFF00`; the runner build uses `/MP8`. The final survival audit peaked at 4.75 GiB process-private and 2,976 MiB DXGI-local memory (about 7.73 GiB when conservatively added). The GPU crossing process peaked at 4.78 GiB private. No run reached the 8 GiB process limit or hit `0xC0000142`.
+
+Evidence logs for this report are beside it: `step6_final3_{cpu,gpu}_{A,B,C}.stdout.log`, `step6_normalpath_near_audit.stdout.log`, `step6_normalpath_near_passage_audit.stdout.log`, `step6_whole_far_packed.stdout.log`, `step6_normalpath_survival_final.stdout.log`, and `step6_normalpath_gpu_crossing_final.stdout.log`. All work and artifacts are under `Standalone\`; no commit, push, or stash was made.

Standalone/Research/OwnerShapeWorld/gpu_field_port/step6_capture_comparison.md

@@ -0,0 +1,16 @@
+# Step 6 CPU/GPU capture comparison
+
+Each pair is a 1280x800 back-buffer capture from the fixed camera. `mae_rgb` and `rmse_rgb` are channel errors on the 0 to 255 scale; pixel percentiles use the maximum RGB channel error per pixel.
+
+| View | MAE RGB | RMSE RGB | p95 pixel max | p99 pixel max | Max channel | Changed pixels | >8 | >32 |
+|---|---:|---:|---:|---:|---:|---:|---:|---:|
+| chamber | 0.6081 | 1.8859 | 3 | 7 | 124 | 42.64% | 0.60% | 0.02% |
+| tunnel | 0.7641 | 3.4371 | 4 | 8 | 130 | 41.72% | 0.96% | 0.28% |
+| passage | 1.6092 | 3.6878 | 12 | 20 | 52 | 50.50% | 8.95% | 0.17% |
+| far | 0.7909 | 2.1558 | 5 | 9 | 128 | 44.83% | 1.25% | 0.02% |
+| passage_floor_closeup | 0.9354 | 2.4817 | 5 | 13 | 100 | 43.14% | 2.54% | 0.05% |
+| tunnel_bend | 0.8236 | 2.1708 | 5 | 11 | 55 | 41.45% | 1.82% | 0.07% |
+| remote_yaw0 | 6.8357 | 19.4185 | 70 | 98 | 125 | 28.35% | 18.65% | 9.85% |
+| remote_yaw90 | 3.3749 | 12.3932 | 36 | 78 | 104 | 18.27% | 9.89% | 5.27% |
+| remote_yaw180 | 4.4919 | 14.7803 | 63 | 78 | 109 | 23.99% | 12.78% | 7.21% |
+| remote_yaw270 | 4.0524 | 13.9541 | 62 | 71 | 126 | 21.05% | 10.69% | 6.45% |

Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl

@@ -27,9 +27,32 @@ cbuffer TileConstants : register(b0)
     float VoxelSizeCm;
     uint OutputBaseVertex;
     uint OutputArgsByteOffset;
-    uint ConstantPad;
+    uint NearFieldMode;
+    uint OutputVertexCapacity;
+    uint3 TileConstantsPadding;
 };
 
+cbuffer FarFieldConstants : register(b1)
+{
+    float FarStrataLo;
+    float FarStrataHi;
+    float FarStrataBlend;
+    uint FarSampleCount;
+    uint FarEvalCounterEnabled;
+    uint FarSpatialBinMask;
+    uint FarSpatialBinSize;
+    uint FarSpatialBinCapacity;
+    uint FarNearDetailEnabled;
+    uint FarFieldSeed;
+    uint FarDetailLayerCount;
+    uint FarFieldConstantPad;
+    int4 FarDetailMinVoxel;
+    int4 FarDetailMaxVoxel;
+};
+
+float NearFieldDensityAt(float3 position, uint clampNearZero,out uint distanceEvaluations);
+StructuredBuffer<float> FarNearNormalDensityInput : register(t18);
+
 cbuffer FarConnectorConstants : register(b2)
 {
     uint FarConnectorCount;
@@ -73,58 +96,74 @@ float LatticeValue(StructuredBuffer<float> density, int3 p)
     return density[SampleIndex(p)];
 }
 
+float GradientLatticeValue(StructuredBuffer<float> density,int3 p)
+{
+    const float cached=LatticeValue(density,p);
+    if(NearFieldMode!=0&&StepVoxels==1&&abs(abs(cached)-4.125f)<1.0e-4f)
+    {
+        // The CPU mesher replaces the fixed proxy-band fallback with the
+        // unclamped analytic field when it samples gradients. Match that
+        // normal path while retaining packed lattice values for classification
+        // and edge positions.
+        return FarNearNormalDensityInput[SampleIndex(p)];
+    }
+    return cached;
+}
+
 float3 LatticeGradient(StructuredBuffer<float> density, int3 p)
 {
     return float3(
-        LatticeValue(density, p + int3(1,0,0)) - LatticeValue(density, p - int3(1,0,0)),
+        GradientLatticeValue(density, p + int3(1,0,0)) - GradientLatticeValue(density, p - int3(1,0,0)),
-        LatticeValue(density, p + int3(0,1,0)) - LatticeValue(density, p - int3(0,1,0)),
+        GradientLatticeValue(density, p + int3(0,1,0)) - GradientLatticeValue(density, p - int3(0,1,0)),
-        LatticeValue(density, p + int3(0,0,1)) - LatticeValue(density, p - int3(0,0,1)));
+        GradientLatticeValue(density, p + int3(0,0,1)) - GradientLatticeValue(density, p - int3(0,0,1)));
 }
 
-float3 EdgePosition(int3 cell, uint edge, StructuredBuffer<float> density,
+float EdgeRootParameter(int3 cell,uint edge,StructuredBuffer<float> density,
-                    StructuredBuffer<uint> midpointIds,
+    StructuredBuffer<uint> midpointIds,StructuredBuffer<MidpointRecord> midpointRecords)
-                    StructuredBuffer<MidpointRecord> midpointRecords)
+{
+    const int2 corners=EdgeCorners[edge];
+    const int3 pa=cell+CornerOffset[corners.x];
+    const int3 pb=cell+CornerOffset[corners.y];
+    const float da=LatticeValue(density,pa),db=LatticeValue(density,pb);
+    const uint midId=midpointIds[CellIndex(uint3(cell))];
+    float mid=0.0f;uint hasMid=0;
+    if(midId!=0xffffffffu)
+    {
+        const MidpointRecord record=midpointRecords[midId];
+        hasMid=(record.metadata.x>>edge)&1;
+        mid=record.values[edge>>2][edge&3];
+    }
+    else if(NearFieldMode!=0&&StepVoxels==1)
+    {
+        const float linearT=VF_FieldRefineRoot(da,db,0.0f,0);
+        if(linearT>=0.375f&&linearT<=0.625f)
+        {
+            const float3 midpoint=(float3(OriginX,OriginY,OriginZ)
+                +float3((pa+pb)*StepVoxels)*0.5f);
+            uint ignoredEvaluations=0;
+            mid=NearFieldDensityAt(midpoint,0,ignoredEvaluations);
+            hasMid=1;
+        }
+    }
+    return VF_FieldRefineRoot(da,db,mid,hasMid);
+}
+
+float3 EdgePosition(int3 cell,uint edge,float t)
 {
     const int2 corners = EdgeCorners[edge];
     const int3 pa = cell + CornerOffset[corners.x];
     const int3 pb = cell + CornerOffset[corners.y];
-    const float da = LatticeValue(density, pa);
-    const float db = LatticeValue(density, pb);
-    const uint midId = midpointIds[CellIndex(uint3(cell))];
-    float mid = 0.0f;
-    uint hasMid = 0;
-    if (midId != 0xffffffffu)
-    {
-        const MidpointRecord record = midpointRecords[midId];
-        hasMid = (record.metadata.x >> edge) & 1;
-        mid = record.values[edge >> 2][edge & 3];
-    }
-    const float t = VF_FieldRefineRoot(da, db, mid, hasMid);
     const float3 originCm = float3(OriginX, OriginY, OriginZ) * VoxelSizeCm;
     const float3 a = originCm + float3(pa * StepVoxels) * VoxelSizeCm;
     const float3 b = originCm + float3(pb * StepVoxels) * VoxelSizeCm;
     return lerp(a, b, t);
 }
 
-float3 EdgeNormal(int3 cell, uint edge, StructuredBuffer<float> density,
+float3 EdgeNormal(int3 cell,uint edge,float t,StructuredBuffer<float> density)
-                  StructuredBuffer<uint> midpointIds,
-                  StructuredBuffer<MidpointRecord> midpointRecords)
 {
     const int2 corners = EdgeCorners[edge];
     const int3 pa = cell + CornerOffset[corners.x];
     const int3 pb = cell + CornerOffset[corners.y];
-    const float da = LatticeValue(density, pa);
-    const float db = LatticeValue(density, pb);
-    const uint midId = midpointIds[CellIndex(uint3(cell))];
-    float mid = 0.0f;
-    uint hasMid = 0;
-    if (midId != 0xffffffffu)
-    {
-        const MidpointRecord record = midpointRecords[midId];
-        hasMid = (record.metadata.x >> edge) & 1;
-        mid = record.values[edge >> 2][edge & 3];
-    }
-    const float t = VF_FieldRefineRoot(da, db, mid, hasMid);
     const float3 ga = LatticeGradient(density, pa);
     const float3 gb = LatticeGradient(density, pb);
     const float3 g = lerp(ga, gb, t);
@@ -145,6 +184,17 @@ struct FarFieldSpatialBin
     uint count;
     uint padding;
 };
+struct FarFieldDetailLayer
+{
+    int lo;
+    int hi;
+    float surfaceRoughness;
+    float roughnessFrequency;
+    float verticalScale;
+    float domainWarpStrength;
+    float domainWarpFrequency;
+    uint roughnessNoiseType;
+};
 struct FarFieldConnectorPanel
 {
     float3 start; float radius;
@@ -171,6 +221,9 @@ StructuredBuffer<uint> FarProtectedAir : register(t13);
 StructuredBuffer<uint> FarProtectedSolid : register(t14);
 RWStructuredBuffer<float> FarFieldDensity : register(u7);
 StructuredBuffer<float> GpuSignDensity : register(t15);
+StructuredBuffer<FarFieldDetailLayer> FarDetailLayers : register(t16);
+StructuredBuffer<uint> FarNearPackedProxy : register(t17);
+RWStructuredBuffer<float> FarNearNormalDensityOutput : register(u5);
 StructuredBuffer<FarFieldConnectorPanel> FarConnectorPanels : register(t8);
 StructuredBuffer<FarConnectorRing> FarConnectorRingInput : register(t9);
 RWStructuredBuffer<FarConnectorRing> FarConnectorRingOutput : register(u7);
@@ -178,18 +231,6 @@ RWByteAddressBuffer FarConnectorDrawArgs : register(u4);
 RWByteAddressBuffer FarConnectorVertexBytes : register(u5);
 RWByteAddressBuffer FarConnectorIndexBytes : register(u6);
 
-cbuffer FarFieldConstants : register(b1)
-{
-    float FarStrataLo;
-    float FarStrataHi;
-    float FarStrataBlend;
-    uint FarSampleCount;
-    uint FarEvalCounterEnabled;
-    uint FarSpatialBinMask;
-    uint FarSpatialBinSize;
-    uint FarSpatialBinCapacity;
-};
-
 RWStructuredBuffer<uint> FarFieldEvalCounters : register(u6);
 
 uint FarFieldBinHash(int3 coord)
@@ -223,6 +264,169 @@ FarFieldSpatialBin FarFieldFindBin(float3 position)
     return emptyBin;
 }
 
+float NearOwnerRoughnessAtStep(float3 position,float surfaceSdf,float caveInfluence,
+    FarFieldDetailLayer layer,uint anchorStep)
+{
+    const float depth=layer.surfaceRoughness*2.0f;
+    if(!(depth>0.0f)||abs(surfaceSdf)>=depth||!(layer.roughnessFrequency>0.0f))return 0.0f;
+    const int step=max(1,(int)anchorStep);
+    const int gridZ=(int)floor(position.z/(float)step)*step;
+    const float sampleZ0=max((float)gridZ,(float)layer.lo);
+    const float sampleZ1=min((float)(gridZ+step),(float)layer.hi);
+    float total;
+    if(sampleZ1>sampleZ0)
+    {
+        const float lower=VF_FieldOwnerRoughnessCore(position.x,position.y,sampleZ0,
+            FarFieldSeed,layer.surfaceRoughness,layer.roughnessFrequency,layer.verticalScale,
+            layer.domainWarpStrength,layer.domainWarpFrequency,layer.roughnessNoiseType);
+        const float upper=VF_FieldOwnerRoughnessCore(position.x,position.y,sampleZ1,
+            FarFieldSeed,layer.surfaceRoughness,layer.roughnessFrequency,layer.verticalScale,
+            layer.domainWarpStrength,layer.domainWarpFrequency,layer.roughnessNoiseType);
+        const float alpha=(position.z-sampleZ0)/(sampleZ1-sampleZ0);
+        total=lower+(upper-lower)*alpha;
+    }
+    else
+        total=VF_FieldOwnerRoughnessCore(position.x,position.y,position.z,
+            FarFieldSeed,layer.surfaceRoughness,layer.roughnessFrequency,layer.verticalScale,
+            layer.domainWarpStrength,layer.domainWarpFrequency,layer.roughnessNoiseType);
+    if(surfaceSdf<0.0f)
+        total=lerp(total,min(total,0.0f),clamp(caveInfluence,0.0f,1.0f));
+    const float fade=1.0f-abs(surfaceSdf)/depth;
+    return total*fade*fade;
+}
+
+float NearOwnerRoughness(float3 position,float surfaceSdf,float caveInfluence,
+    float passageInfluence,FarFieldDetailLayer layer)
+{
+    const float passage=saturate(passageInfluence);
+    const float anchored=NearOwnerRoughnessAtStep(position,surfaceSdf,caveInfluence,layer,8);
+    if(passage<=0.40f)return anchored;
+    const float direct=NearOwnerRoughnessAtStep(position,surfaceSdf,caveInfluence,layer,2);
+    if(passage>=0.60f)return direct;
+    const float directWeight=VF_FieldSmoothUnit((passage-0.40f)/0.20f);
+    const float anchoredWeight=1.0f-directWeight;
+    const float normalization=1.0f/sqrt(anchoredWeight*anchoredWeight+directWeight*directWeight);
+    return (anchoredWeight*anchored+directWeight*direct)*normalization;
+}
+
+float NearFieldDensityAt(float3 position,uint clampNearZero,out uint distanceEvaluations)
+{
+    const FarFieldSpatialBin range=FarFieldFindBin(position);
+    float cave=100000.0f,overlay=-100000.0f;
+    float caveFloorInfluence=0.0f,cavePassageInfluence=0.0f;
+    distanceEvaluations=0;
+    [loop]
+    for(uint item=0;item<range.count;++item)
+    {
+        const VF_FieldPrimitive shape=FarShapes[FarSpatialBinShapeIds[range.offset+item]];
+        if((shape.flags&VF_FIELD_PRIMITIVE_CARVE)!=0)
+        {
+            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;
+            float distance=VF_FieldPrimitiveDistance(shape,position.x,position.y,position.z);
+            ++distanceEvaluations;
+            float primitiveFloorInfluence=0.0f;
+            if((shape.flags&VF_FIELD_PRIMITIVE_HAS_FLOOR)!=0)
+            {
+                const float floor=VF_FieldPrimitiveFloorAt(shape,position.x,position.y);
+                const float floorT=(abs(position.z-floor)-shape.floorBlendRadius)/12.0f;
+                primitiveFloorInfluence=1.0f-VF_FieldSmoothUnit(floorT);
+                distance=VF_FieldSmoothMax(distance,floor-position.z,shape.floorBlendRadius);
+            }
+            const float metadataWidth=max(shape.blendRadius,6.0f);
+            const float primitiveWeight=VF_FieldSmoothUnit(0.5f+(cave-distance)/(2.0f*metadataWidth));
+            caveFloorInfluence=lerp(caveFloorInfluence,primitiveFloorInfluence,primitiveWeight);
+            if(cavePassageInfluence>0.0f||(shape.flags&VF_FIELD_PRIMITIVE_PASSAGE)!=0)
+                cavePassageInfluence=lerp(cavePassageInfluence,
+                    (shape.flags&VF_FIELD_PRIMITIVE_PASSAGE)!=0?1.0f:0.0f,primitiveWeight);
+            cave=VF_FieldSmoothMin(cave,distance,shape.blendRadius);
+        }
+        else if(position.x>=shape.minVoxel[0]&&position.x<=shape.maxVoxel[0]
+            &&position.y>=shape.minVoxel[1]&&position.y<=shape.maxVoxel[1]
+            &&position.z>=shape.minVoxel[2]&&position.z<=shape.maxVoxel[2])
+        {
+            overlay=max(overlay,-VF_FieldPrimitiveDistance(shape,position.x,position.y,position.z));
+            ++distanceEvaluations;
+        }
+    }
+    const float base=min(position.z-FarStrataLo,FarStrataHi-position.z);
+    const float carveOrOverlay=max(cave,overlay);
+    const float caveVsOverlay=VF_FieldSmoothUnit(0.5f+(cave-overlay)/12.0f);
+    const float carveInfluence=VF_FieldSmoothUnit(0.5f+(base-carveOrOverlay)/12.0f);
+    const float caveInfluence=carveInfluence*caveVsOverlay;
+    const float floorInfluence=caveInfluence*caveFloorInfluence;
+    const float passageInfluence=caveInfluence*cavePassageInfluence;
+    const float surfaceSdf=base<=carveOrOverlay?base:(cave>=overlay?cave:overlay);
+    float rockDistance=VF_FieldSmoothMin(base,carveOrOverlay,FarStrataBlend);
+    if(NearFieldMode!=0
+        &&all(position>=float3(FarDetailMinVoxel.xyz))
+        &&all(position<=float3(FarDetailMaxVoxel.xyz))
+        &&abs(rockDistance)<=4.125f)
+    {
+        const float macro=VF_FieldValueNoise(position.x,position.y,position.z,32,
+            FarFieldSeed^0xB5297A4Du);
+        const float mid=VF_FieldValueNoise(position.x,position.y,position.z,16,
+            FarFieldSeed^0x63D83595u);
+        const float fine=VF_FieldValueNoise(position.x,position.y,position.z,4,
+            FarFieldSeed^0xA511E9B3u);
+        const float relief=1.20f*macro+(0.80f+0.25f*macro)*(0.40f*mid+0.60f*fine);
+        const float passageBandRelief=1.20f*macro+(0.80f+0.25f*macro)*(0.40f*mid);
+        float ownerRelief=0.0f;
+        [loop]
+        for(uint layerIndex=0;layerIndex<FarDetailLayerCount;++layerIndex)
+        {
+            const FarFieldDetailLayer layer=FarDetailLayers[layerIndex];
+            if(position.z>=layer.lo&&position.z<=layer.hi)
+            { ownerRelief=NearOwnerRoughness(position,surfaceSdf,caveInfluence,
+                passageInfluence,layer);break; }
+        }
+        const float caveFactor=saturate(caveInfluence);
+        const float floorFactor=min(caveFactor,saturate(floorInfluence));
+        const float passageFactor=saturate(passageInfluence);
+        const float nonPassage=1.0f-passageFactor;
+        const float strataWeight=nonPassage*(1.0f-caveFactor);
+        const float caveWeight=nonPassage*(caveFactor-floorFactor);
+        const float floorWeight=nonPassage*floorFactor;
+        const float detail=strataWeight*(1.35f*(relief+0.5f*ownerRelief))
+            +caveWeight*relief+floorWeight*(0.55f*relief)
+            +passageFactor*(0.75f*passageBandRelief);
+
+        float ownerBound=0.0f;
+        [loop]
+        for(uint layerIndex=0;layerIndex<FarDetailLayerCount;++layerIndex)
+        {
+            const FarFieldDetailLayer layer=FarDetailLayers[layerIndex];
+            if(position.z>=layer.lo&&position.z<=layer.hi)
+            {
+                const float depth=layer.surfaceRoughness*2.0f;
+                if(depth>0.0f&&abs(surfaceSdf)<depth&&layer.roughnessFrequency>0.0f)
+                {
+                    const float fade=1.0f-abs(surfaceSdf)/depth;
+                    ownerBound=3.15f*1.25f*layer.surfaceRoughness*fade*fade;
+                }
+                break;
+            }
+        }
+        const float reliefBound=1.20f*abs(macro)+(0.80f+0.25f*macro);
+        const float contrast=1.35f*(1.0f-caveFactor)+1.00f*(caveFactor-floorFactor)+0.55f*floorFactor;
+        const float localSupremum=min(3.9f,(reliefBound+0.5f*ownerBound)*contrast);
+        const float edge=clamp(localSupremum+0.125f,1.5f,4.125f);
+        const float distance=abs(rockDistance);
+        float fade=1.0f;
+        if(distance>1.5f)
+        {
+            const float width=edge-1.5f;
+            fade=width>0.0f?VF_FieldSmoothUnit((edge-distance)/width):0.0f;
+        }
+        rockDistance+=clamp(detail,-3.9f,3.9f)*fade;
+    }
+    float density=-rockDistance;
+    if(clampNearZero!=0&&abs(density)<0.05f)
+        density=rockDistance>=0.0f?-0.05f:0.05f;
+    return density;
+}
+
 [numthreads(64,1,1)]
 void EvaluateFarFieldCS(uint3 dispatchId : SV_DispatchThreadID)
 {
@@ -233,50 +437,72 @@ void EvaluateFarFieldCS(uint3 dispatchId : SV_DispatchThreadID)
     const uint z = index / (SampleDim * SampleDim);
     const float3 position = float3(OriginX, OriginY, OriginZ)
         + float3(int(x) - 1, int(y) - 1, int(z) - 1) * StepVoxels;
-    const FarFieldSpatialBin range = FarFieldFindBin(position);
+    float density;
-    float cave = 100000.0f;
+    uint distanceEvaluations=0;
-    float overlay = -100000.0f;
+    float nearNormalDensity=0.0f;
-    uint distanceEvaluations = 0;
+    if(NearFieldMode!=0)
-    [loop]
-    for (uint item = 0; item < range.count; ++item)
     {
-        const VF_FieldPrimitive shape = FarShapes[FarSpatialBinShapeIds[range.offset + item]];
+        nearNormalDensity=NearFieldDensityAt(position,0,distanceEvaluations);
-        if ((shape.flags & VF_FIELD_PRIMITIVE_CARVE) != 0)
+        density=nearNormalDensity;
+        if(abs(density)<0.05f)
+            density=density<0.0f?-0.05f:(density>0.0f?0.05f:-0.05f);
+        const uint packedProxy=FarNearPackedProxy[index];
+        if(packedProxy!=0xffffffffu)
         {
-            if (position.x < shape.minVoxel[0] - 6 || position.x > shape.maxVoxel[0] + 6
+            // The CPU uploads sparse interpolation magnitudes. The GPU decides
-                || position.y < shape.minVoxel[1] - 6 || position.y > shape.maxVoxel[1] + 6
+            // ordinary signs, except for the one-quantum band where float sign
-                || position.z < shape.minVoxel[2] - 6 || position.z > shape.maxVoxel[2] + 6)
+            // noise can flip an entire cell against the packed CPU cache.
-                continue;
+            const bool packedSignOverride=(packedProxy&0x40000000u)!=0u;
-            float distance = VF_FieldPrimitiveDistance(shape,
+            const bool packedRockSign=(packedProxy&0x80000000u)!=0u;
-                position.x, position.y, position.z);
+            const uint magnitudeQ=packedProxy&0x3fffffffu;
-            ++distanceEvaluations;
+            const float magnitude=magnitudeQ==0u?4.125f:(float)magnitudeQ*(1.0f/256.0f);
-            if ((shape.flags & VF_FIELD_PRIMITIVE_HAS_FLOOR) != 0)
+            const bool negative=packedSignOverride?packedRockSign:(density<=0.0f);
-            {
+            density=negative?-magnitude:magnitude;
-                const float floorZ = VF_FieldPrimitiveFloorAt(shape, position.x, position.y);
-                distance = VF_FieldSmoothMax(distance, floorZ - position.z,
-                    shape.floorBlendRadius);
-            }
-            cave = VF_FieldSmoothMin(cave, distance, shape.blendRadius);
-        }
-        else if (position.x >= shape.minVoxel[0] && position.x <= shape.maxVoxel[0]
-            && position.y >= shape.minVoxel[1] && position.y <= shape.maxVoxel[1]
-            && position.z >= shape.minVoxel[2] && position.z <= shape.maxVoxel[2])
-        {
-            overlay = max(overlay, -VF_FieldPrimitiveDistance(shape,
-                position.x, position.y, position.z));
-            ++distanceEvaluations;
         }
     }
-    const float base = min(position.z - FarStrataLo, FarStrataHi - position.z);
+    else
-    float rockDistance = VF_FieldSmoothMin(base, max(cave, overlay), FarStrataBlend);
+    {
+        const FarFieldSpatialBin range=FarFieldFindBin(position);
+        float cave=100000.0f,overlay=-100000.0f;
+        [loop]
+        for(uint item=0;item<range.count;++item)
+        {
+            const VF_FieldPrimitive shape=FarShapes[FarSpatialBinShapeIds[range.offset+item]];
+            if((shape.flags&VF_FIELD_PRIMITIVE_CARVE)!=0)
+            {
+                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;
+                float distance=VF_FieldPrimitiveDistance(shape,position.x,position.y,position.z);
+                ++distanceEvaluations;
+                if((shape.flags&VF_FIELD_PRIMITIVE_HAS_FLOOR)!=0)
+                {
+                    const float floorZ=VF_FieldPrimitiveFloorAt(shape,position.x,position.y);
+                    distance=VF_FieldSmoothMax(distance,floorZ-position.z,shape.floorBlendRadius);
+                }
+                cave=VF_FieldSmoothMin(cave,distance,shape.blendRadius);
+            }
+            else if(position.x>=shape.minVoxel[0]&&position.x<=shape.maxVoxel[0]
+                &&position.y>=shape.minVoxel[1]&&position.y<=shape.maxVoxel[1]
+                &&position.z>=shape.minVoxel[2]&&position.z<=shape.maxVoxel[2])
+            {
+                overlay=max(overlay,-VF_FieldPrimitiveDistance(shape,position.x,position.y,position.z));
+                ++distanceEvaluations;
+            }
+        }
+        const float base=min(position.z-FarStrataLo,FarStrataHi-position.z);
+        const float rockDistance=VF_FieldSmoothMin(base,max(cave,overlay),FarStrataBlend);
+        const float rawDensity=-rockDistance;
+        const float magnitude=max(abs(rawDensity),0.05f);
+        density=rockDistance>=0.0f?-magnitude:magnitude;
+    }
     const uint protectedAir = (FarProtectedAir[index >> 5] >> (index & 31)) & 1;
     const uint protectedSolid = (FarProtectedSolid[index >> 5] >> (index & 31)) & 1;
-    const float rawDensity = -rockDistance;
+    const float protectedMinimum=NearFieldMode!=0?1.0f/256.0f:0.05f;
-    const float magnitude = max(abs(rawDensity), 0.05f);
+    if(protectedAir!=0)density=max(abs(density),protectedMinimum);
-    float density = rockDistance >= 0.0f ? -magnitude : magnitude;
+    else if(protectedSolid!=0)density=-max(abs(density),protectedMinimum);
-    if (protectedAir != 0) density = magnitude;
-    else if (protectedSolid != 0) density = -magnitude;
     FarFieldDensity[index] = density;
+    FarNearNormalDensityOutput[index]=nearNormalDensity;
     if (FarEvalCounterEnabled != 0) FarFieldEvalCounters[index] = distanceEvaluations;
 }
 
@@ -320,12 +546,12 @@ float FarRockDistanceAt(float3 position,int3 cachedBinCoord,FarFieldSpatialBin c
 
 float FarConnectorRadius(float3 center, float3 direction, float nominalRadius)
 {
-    // Keep the bounded signed-distance refinement, but stop once the residual
-    // reaches the same acceptance bound used by the original fallback. A
-    // spatial-bin cache avoids re-hashing nearby probes.
     const float3 nominalPosition=center+direction*nominalRadius;
     const int3 cachedBinCoord=int3(floor(nominalPosition/max(1.0f,(float)FarSpatialBinSize)));
     const FarFieldSpatialBin cachedRange=FarFieldFindBin(nominalPosition);
+    // Keep the bounded signed-distance refinement, but stop once the residual
+    // reaches the same acceptance bound used by the original fallback. A
+    // spatial-bin cache avoids re-hashing nearby probes.
     const float maxOffset=max(16.0f,nominalRadius*0.75f);
     const float minRadius=max(0.0f,nominalRadius-maxOffset);
     const float maxRadius=nominalRadius+maxOffset;
@@ -470,8 +696,8 @@ void BuildFarConnectorRingsCS(uint3 groupId : SV_GroupID,
         }
         if (ringId == 0)
         {
-            const uint maxVerticesPerTile = CellCount * 15;
+            const uint slot = OutputVertexCapacity > 0
-            const uint slot = maxVerticesPerTile > 0 ? OutputBaseVertex / maxVerticesPerTile : 0;
+                ? OutputBaseVertex / OutputVertexCapacity : 0;
             const uint panelBaseVertex = slot * FarConnectorPanelVertexCapacity;
             FarConnectorDrawArgs.Store(OutputArgsByteOffset + 20, FarConnectorVertexCount);
             FarConnectorDrawArgs.Store(OutputArgsByteOffset + 24, FarConnectorVertexCount > 0 ? 1 : 0);
@@ -547,8 +773,8 @@ 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 maxVerticesPerTile = CellCount * 15;
+    const uint slot = OutputVertexCapacity > 0
-    const uint slot = maxVerticesPerTile > 0 ? OutputBaseVertex / maxVerticesPerTile : 0;
+        ? OutputBaseVertex / OutputVertexCapacity : 0;
     const uint base = slot * FarConnectorPanelVertexCapacity;
     const uint byteAddress = (base + vertexId) * 32;
     FarConnectorVertexBytes.Store3(byteAddress, asuint(vertex.position));
@@ -644,7 +870,7 @@ void ScanGroupsCS(uint3 dispatchId : SV_DispatchThreadID)
         GroupOffsets[i] = offset;
         offset += InputGroupVertexCounts[i];
     }
-    DrawArgs.Store(OutputArgsByteOffset + 0, offset);
+    DrawArgs.Store(OutputArgsByteOffset + 0, min(offset, OutputVertexCapacity));
     DrawArgs.Store(OutputArgsByteOffset + 4, 1);
     DrawArgs.Store(OutputArgsByteOffset + 8, OutputBaseVertex);
     DrawArgs.Store(OutputArgsByteOffset + 12, 0);
@@ -672,11 +898,13 @@ void EmitCS(uint3 dispatchId : SV_DispatchThreadID)
     const uint index = dispatchId.x;
     if (index >= CellCount) return;
     const uint caseIndex = Cases[index];
-    const uint base = OutputBaseVertex + LocalOffsets[index] + PrefixGroups[index / 256];
+    const uint localBase = LocalOffsets[index] + PrefixGroups[index / 256];
+    if(localBase>=OutputVertexCapacity)return;
+    const uint base = OutputBaseVertex + localBase;
     const uint3 cell = uint3(index % CellsPerAxis,
         (index / CellsPerAxis) % CellsPerAxis,
         index / (CellsPerAxis * CellsPerAxis));
-    [unroll]
+    [loop]
     for (uint slot = 0; slot < 15; ++slot)
     {
         uint tableSlot = slot;
@@ -684,10 +912,12 @@ void EmitCS(uint3 dispatchId : SV_DispatchThreadID)
         else if (slot % 3 == 2) tableSlot = slot - 1;
         const int edgeValue = Triangles[caseIndex * 16 + tableSlot];
         if (edgeValue < 0) break;
+        if(localBase+slot>=OutputVertexCapacity)break;
         const uint edge = uint(edgeValue);
+        const float rootT=EdgeRootParameter(int3(cell),edge,Density,MidpointCellIds,MidpointData);
         GpuVertex v;
-        v.position = EdgePosition(int3(cell), edge, Density, MidpointCellIds, MidpointData);
+        v.position = EdgePosition(int3(cell),edge,rootT);
-        v.normal = EdgeNormal(int3(cell), edge, Density, MidpointCellIds, MidpointData);
+        v.normal = EdgeNormal(int3(cell),edge,rootT,Density);
         v.uv = v.position.xy / VoxelSizeCm;
         const uint byteAddress = (base + slot) * 32;
         VertexBytes.Store3(byteAddress, asuint(v.position));

Standalone/Research/OwnerShapeWorld/speed_round/round9/REPORT.md

@@ -69,7 +69,16 @@ The bad-shell control fails as required: it removes 512 shell triangles and dete
 
 The six fixed views from the final executable are pixel-identical to Round 8. The [comparison](capture_compare_final.log) confirms zero changed pixels in all six. Images are in this folder: chamber, tunnel, passage, far, passage floor close-up, and tunnel bend.
 
-The first remote capture attempt used an assumed camera position of `(500, -123.25, -19.85)` m and is invalid: it does not match Round 8’s X=500 views. I asked for the original Y/Z coordinates and will replace those four provisional images and update this section when they are provided. The captures do not alter the mesh hashes or the fixed-view result.
+The remote comparison was recaptured with the recorded camera positions `(500,0,0,yaw,0)` for yaw 0/90/180/270. Round 8 (`VoxelForgeShapeRound8Final.exe`) and Round 9 (`VoxelForgeRound9Final.exe`) are exact at yaw 0. The remaining differences are at single-channel precision except for six channel values at yaw 270:
+
+| View | Changed pixels | Changed channels | Mean channel delta | Max channel delta |
+| --- | ---: | ---: | ---: | ---: |
+| Remote yaw 0 | 0 / 1,024,000 | 0 / 3,072,000 | 0.000000 | 0 |
+| Remote yaw 90 | 111 / 1,024,000 | 154 / 3,072,000 | 0.000050 | 1 |
+| Remote yaw 180 | 32 / 1,024,000 | 44 / 3,072,000 | 0.000014 | 1 |
+| Remote yaw 270 | 386 / 1,024,000 | 1,040 / 3,072,000 | 0.000695 | 6 |
+
+Both executables were run fresh with `--capture-after-clipmap --shape-world --force-generate --workers=8 --clip-radius=3`; all eight capture processes exited 0 under BelowNormal priority and affinity `0xFF00`. Images and the [pixel comparison](../round10/remote_compare.log) are in `../round10/`.
 
 ## Failed builds and runs
 

Standalone/Source/StandaloneViewer.cpp

@@ -1372,6 +1372,7 @@ namespace
         uint32 TileCount=0;
         uint64 TimestampFrequency=0;
         double GpuDispatchMs=0.0;
+        std::chrono::steady_clock::time_point SubmittedAt{};
         bool bGpuComplete=false,bTimingResolved=false;
     };
 
@@ -1502,6 +1503,14 @@ namespace
         std::unordered_map<FVoxelTileKey,VoxelForgeCore::FarFieldTileInput,FStandaloneTileKeyHash> GpuFarPreparedInputs;
         std::unordered_map<FVoxelTileKey,std::string,FStandaloneTileKeyHash> GpuFarPreparationErrors;
         double GpuFarPreparationWallMs=0.0,GpuFarPreparationWorkerMs=0.0;
+        uint64 GpuFarPumpCalls=0,GpuFarCompletionPolls=0;
+        double GpuFarFirstPumpMs=-1.0,GpuFarLastPumpMs=-1.0;
+        double GpuFarFirstSubmitMs=-1.0,GpuFarLastSubmitMs=-1.0;
+        double GpuFarFirstCompletionPollMs=-1.0,GpuFarLastCompletionPollMs=-1.0;
+        double GpuFarPumpWallMs=0.0,GpuFarPumpMaxWallMs=0.0;
+        double GpuFarPumpGapMs=0.0,GpuFarPumpGapMaxMs=0.0;
+        std::chrono::steady_clock::time_point LastGpuFarPumpAt{};
+        std::chrono::steady_clock::time_point LastDrawEndAt{};
         VoxelForgeCore::DetailVariant ShapeDetailVariant=VoxelForgeCore::DetailVariant::coherent;
         bool bWireframe=false;
         bool bCaptured=false;
@@ -1610,6 +1619,12 @@ namespace
             GpuFarFreeOutputSlots.clear();GpuFarNextOutputSlot=0;
             GpuFarPreparedInputs.clear();GpuFarPreparationErrors.clear();
             GpuFarPreparationWallMs=0.0;GpuFarPreparationWorkerMs=0.0;
+            GpuFarPumpCalls=0;GpuFarCompletionPolls=0;
+            GpuFarFirstPumpMs=GpuFarLastPumpMs=GpuFarFirstSubmitMs=GpuFarLastSubmitMs=-1.0;
+            GpuFarFirstCompletionPollMs=GpuFarLastCompletionPollMs=-1.0;
+            GpuFarPumpWallMs=GpuFarPumpMaxWallMs=GpuFarPumpGapMs=GpuFarPumpGapMaxMs=0.0;
+            LastGpuFarPumpAt={};
+            LastDrawEndAt={};
             GpuFarAudit.reset();bGpuFarAuditReported=false;
             GpuFarTilesSubmitted=0;GpuFarEnqueueWallSeconds=0.0;GpuFarGpuDispatchMs=0.0;
             GpuFarGpuDispatchMaxMs=0.0;GpuFarSubmissionSpanMs=0.0;GpuFarSubmissionSpanMaxMs=0.0;
@@ -1916,8 +1931,15 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             Request.includeFarWorld=bIncludeFarWorld;Request.roundTunnels=bRoundTunnels;
             Request.enhancedTunnelWalls=bRoundTunnelEnhanced;
             Request.profileDetailBand=bProfileDetailBand;
+            const auto ShapeCompileStart=std::chrono::steady_clock::now();
+            std::printf("[ColdTimeline] stage=shape_compile begin_ms=%.3f\n",
+                std::chrono::duration<double,std::milli>(ShapeCompileStart-ColdProcessStart).count());
             if(!ShapeWorld->compile(Request,CoreScheduler,Error))
             {ShapeWorld.reset();return false;}
+            const auto ShapeCompileEnd=std::chrono::steady_clock::now();
+            std::printf("[ColdTimeline] stage=shape_compile end_ms=%.3f duration_ms=%.3f\n",
+                std::chrono::duration<double,std::milli>(ShapeCompileEnd-ColdProcessStart).count(),
+                std::chrono::duration<double,std::milli>(ShapeCompileEnd-ShapeCompileStart).count());
             bShapePipelineActive=bNeighborhoodOnly;
             bNeighborhoodReadyReported=false;
             bFarShapeStagePending=bShapeWorld&&!bNeighborhoodOnly&&!ShapeWorld->hasFarWorld();
@@ -2014,9 +2036,13 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             StopGeneration();
             ReleaseMeshes();
             GenerationStart=std::chrono::steady_clock::now();
+            std::printf("[ColdTimeline] stage=generation_setup begin_ms=%.3f\n",
+                std::chrono::duration<double,std::milli>(GenerationStart-ColdProcessStart).count());
             std::string Error;
             if(!ConfigureWorld(Error))
             {std::fprintf(stderr,"World load failed: %s\n",Error.c_str());MessageBoxA(Window,Error.c_str(),"VoxelForge viewer",MB_ICONERROR);return;}
+            std::printf("[ColdTimeline] stage=generation_setup configured_ms=%.3f\n",
+                std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count());
             const FVector CameraVoxels(Position.X*4.0,Position.Y*4.0,Position.Z*4.0);
             BuildTargetChunk=StandaloneWorldMesher::CameraChunk(CameraVoxels);
             bFullDetailReported=false;bInitialClipmapReported=false;bAutoCaptureWritten=false;
@@ -2085,7 +2111,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 {
                     std::vector<FVoxelTileKey> AuditKeys;AuditKeys.reserve(GpuFarEligibleTiles.size());
                     for(const FVoxelTileKey& Key:DesiredTiles)
-                        if(GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key);
+                        if(Key.Level>=6&&GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key);
                     GpuFarAudit=std::make_unique<FFarSurvivalAudit>(*World,*ShapeWorld,AuditKeys,false);
                 }
             }
@@ -2104,6 +2130,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             InitialWork.SubmittedJobs=Update.NewJobs;InitialWork.PendingJobs=Update.NewJobs;
             InitialWork.StartTime=GenerationStart;
             bGenerationRunning=true;bRebuildPending=false;
+            std::printf("[ColdTimeline] stage=cpu_stream_start submitted_ms=%.3f cpu_jobs=%d gpu_queued=%zu\n",
+                std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count(),
+                Update.NewJobs,GpuFarQueue.size());
             const char* TunnelMode=!bRoundTunnels?"":(bRoundTunnelEnhanced
                 ?" with detailed round/winding tunnels":" with legacy round/winding tunnels");
             std::printf("Streaming seed %d%s%s: %d desired tiles, %d full-detail tiles, %d workers; upload budget %d batches / %.1f ms per frame.\n",
@@ -2147,7 +2176,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 {
                     std::vector<FVoxelTileKey> AuditKeys;AuditKeys.reserve(GpuFarEligibleTiles.size());
                     for(const FVoxelTileKey& Key:DesiredTiles)
-                        if(GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key);
+                        if(Key.Level>=6&&GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key);
                     GpuFarAudit=std::make_unique<FFarSurvivalAudit>(*World,*ShapeWorld,AuditKeys,false);
                 }
             }
@@ -2163,6 +2192,18 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             std::fflush(stdout);
         }
 
+        void RecordGpuCrossingTileReady(const FVoxelTileKey& Key)
+        {
+            const auto ReadyAt=std::chrono::steady_clock::now();
+            for(auto& Pair:CrossingWork)
+            {
+                FCrossingWork& Work=Pair.second;
+                if(Work.bHasNearestNewTile&&Work.NearestNewTile==Key
+                    &&Work.NearestReadySeconds<0.0&&Work.StartTime.time_since_epoch().count()!=0)
+                    Work.NearestReadySeconds=std::chrono::duration<double>(ReadyAt-Work.StartTime).count();
+            }
+        }
+
         bool IsRenderedTileReady(const FVoxelTileKey& Key)
         {
             if(bGpuFarEnabled&&ShapeWorld&&ShapeWorld->hasFarWorld()
@@ -2178,6 +2219,10 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                         &&Context->GetData(Batch.CompletionEvent,nullptr,0,D3D11_ASYNC_GETDATA_DONOTFLUSH)==S_OK)
                     {
                         Batch.bGpuComplete=true;
+                        const auto PollAt=std::chrono::steady_clock::now();
+                        const double PollMs=std::chrono::duration<double,std::milli>(PollAt-ColdProcessStart).count();
+                        if(GpuFarFirstCompletionPollMs<0.0)GpuFarFirstCompletionPollMs=PollMs;
+                        GpuFarLastCompletionPollMs=PollMs;++GpuFarCompletionPolls;
                         Batch.CompletionEvent->Release();Batch.CompletionEvent=nullptr;
                     }
                     if(Batch.bGpuComplete&&!Batch.bTimingResolved)
@@ -2202,6 +2247,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                             GpuFarGpuDispatchMaxMs=std::max(GpuFarGpuDispatchMaxMs,TileDispatchMs);
                         }
                     }
+                    if(Batch.bGpuComplete)RecordGpuCrossingTileReady(Key);
                     return Batch.bGpuComplete;
                 }
                 if(!Tile.bGpuComplete&&Tile.CompletionEvent&&Context
@@ -2224,6 +2270,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                         Tile.bTimingResolved=true;
                     }
                 }
+                if(Tile.bGpuComplete)RecordGpuCrossingTileReady(Key);
                 return Tile.bGpuComplete;
             }
             return ResidentTiles.contains(Key);
@@ -2233,6 +2280,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
         {
             if(NeedPreparation.empty())return;
             const auto PrepareStart=std::chrono::steady_clock::now();
+            std::printf("[ColdTimeline] stage=gpu_input_export begin_ms=%.3f tiles=%zu\n",
+                std::chrono::duration<double,std::milli>(PrepareStart-ColdProcessStart).count(),NeedPreparation.size());
             std::vector<std::optional<VoxelForgeCore::FarFieldTileInput>> Inputs(NeedPreparation.size());
             std::vector<std::string> Errors(NeedPreparation.size());
             const std::size_t ThreadCount=std::min<std::size_t>(8,NeedPreparation.size());
@@ -2250,7 +2299,12 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                         if(Index>=NeedPreparation.size())break;
                         const FVoxelTileKey& Key=NeedPreparation[Index];
                         const VoxelForgeCore::TileKey CoreKey{{Key.Coord.X,Key.Coord.Y,Key.Coord.Z},Key.Level};
-                        try { Inputs[Index]=ShapeWorld->exportFarFieldTile(CoreKey,Errors[Index]); }
+                        try
+                        {
+                            Inputs[Index]=Key.Level<World->FullResClipLevels
+                                ?ShapeWorld->exportGpuFieldTile(CoreKey,Errors[Index])
+                                :ShapeWorld->exportFarFieldTile(CoreKey,Errors[Index]);
+                        }
                         catch(const std::exception& Error) { Errors[Index]=Error.what(); }
                     }
                     WorkerElapsed[Worker]=std::chrono::duration<double,std::milli>(
@@ -2275,6 +2329,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             }
             std::printf("[GpuFarInputPrepare] tiles=%zu workers=%zu wall_ms=%.3f worker_sum_ms=%.3f failed=%zu\n",
                 NeedPreparation.size(),ThreadCount,PreparationMs,WorkerSumMs,Failed);
+            std::printf("[ColdTimeline] stage=gpu_input_export end_ms=%.3f duration_ms=%.3f worker_sum_ms=%.3f\n",
+                std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count(),
+                PreparationMs,WorkerSumMs);
             std::fflush(stdout);
         }
 
@@ -2287,9 +2344,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             std::vector<FVoxelTileKey> Eligible;Eligible.reserve(Keys.size());
             for(const FVoxelTileKey& Key:Keys)
             {
-                if(Key.Level<World->FullResClipLevels){CpuTiles.push_back(Key);continue;}
                 const VoxelForgeCore::TileKey CoreKey{{Key.Coord.X,Key.Coord.Y,Key.Coord.Z},Key.Level};
-                if(ShapeWorld->isFarFieldTileEligible(CoreKey))
+                const bool NearGpuEligible=Key.Level<World->FullResClipLevels;
+                if(NearGpuEligible||ShapeWorld->isFarFieldTileEligible(CoreKey))
                 {
                     Eligible.push_back(Key);
                 }
@@ -2306,6 +2363,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             if(!Eligible.empty())
             {
                 std::string PoolError;const auto ReserveStart=std::chrono::steady_clock::now();
+                std::printf("[ColdTimeline] stage=gpu_output_pool begin_ms=%.3f slots=%u\n",
+                    std::chrono::duration<double,std::milli>(ReserveStart-ColdProcessStart).count(),
+                    GpuFarOutputCapacitySlots);
                 if(!VoxelForgeGpuFarViewer::ReserveFarOutputPool(Device,Context,GpuFarShaderPath,
                     GpuFarOutputCapacitySlots,static_cast<uint32>(World->CoarseTileCells),PoolError))
                 {
@@ -2315,6 +2375,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 }
                 GpuFarPoolReserveMs+=std::chrono::duration<double,std::milli>(
                     std::chrono::steady_clock::now()-ReserveStart).count();
+                std::printf("[ColdTimeline] stage=gpu_output_pool end_ms=%.3f duration_ms=%.3f\n",
+                    std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count(),
+                    GpuFarPoolReserveMs);
                 SampleGpuFarMemory();
             }
             std::vector<FVoxelTileKey> NeedPreparation;NeedPreparation.reserve(Eligible.size());
@@ -2322,24 +2385,58 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 if(!GpuFarTiles.contains(Key))NeedPreparation.push_back(Key);
             if(!NeedPreparation.empty())
                 PrepareGpuFarInputs(std::move(NeedPreparation));
+            std::vector<FVoxelTileKey> PreparedEligible;PreparedEligible.reserve(Eligible.size());
             for(const FVoxelTileKey& Key:Eligible)
             {
+                if(!GpuFarTiles.contains(Key)&&!GpuFarPreparedInputs.contains(Key))
+                {
+                    CpuTiles.push_back(Key);
+                    const auto Error=GpuFarPreparationErrors.find(Key);
+                    std::fprintf(stderr,"[GpuFieldInputFallback] key=(%d,%d,%d,lod%d) reason=%s\n",
+                        Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,
+                        Error==GpuFarPreparationErrors.end()?"missing prepared input":Error->second.c_str());
+                    continue;
+                }
                 GpuFarEligibleTiles.insert(Key);
                 if(!GpuFarTiles.contains(Key))
                 {GpuFarQueue.push_back(Key);GpuFarQueued.insert(Key);}
+                PreparedEligible.push_back(Key);
             }
+            Eligible.swap(PreparedEligible);
+            std::sort(GpuFarQueue.begin(),GpuFarQueue.end(),[&](const FVoxelTileKey& A,const FVoxelTileKey& B)
+            {
+                if(A.Level!=B.Level)return A.Level<B.Level;
+                const FIntVector AO=A.OriginVoxels(),BO=B.OriginVoxels();
+                const double AC=(AO.X+0.5*A.ExtentVoxels()-Position.X*4.0)*(AO.X+0.5*A.ExtentVoxels()-Position.X*4.0)
+                    +(AO.Y+0.5*A.ExtentVoxels()-Position.Y*4.0)*(AO.Y+0.5*A.ExtentVoxels()-Position.Y*4.0)
+                    +(AO.Z+0.5*A.ExtentVoxels()-Position.Z*4.0)*(AO.Z+0.5*A.ExtentVoxels()-Position.Z*4.0);
+                const double BC=(BO.X+0.5*B.ExtentVoxels()-Position.X*4.0)*(BO.X+0.5*B.ExtentVoxels()-Position.X*4.0)
+                    +(BO.Y+0.5*B.ExtentVoxels()-Position.Y*4.0)*(BO.Y+0.5*B.ExtentVoxels()-Position.Y*4.0)
+                    +(BO.Z+0.5*B.ExtentVoxels()-Position.Z*4.0)*(BO.Z+0.5*B.ExtentVoxels()-Position.Z*4.0);
+                return AC<BC;
+            });
+            GpuFarQueued.clear();for(const FVoxelTileKey& Key:GpuFarQueue)GpuFarQueued.insert(Key);
             if(Streamer)Streamer->SetDesiredTileFilter(CpuTiles);
-            const std::size_t CpuTransitionFallbacks=static_cast<std::size_t>(std::count_if(
+            const std::size_t NearGpuTiles=static_cast<std::size_t>(std::count_if(
-                Keys.begin(),Keys.end(),[&](const FVoxelTileKey& Key)
+                GpuFarEligibleTiles.begin(),GpuFarEligibleTiles.end(),[&](const FVoxelTileKey& Key)
-                    {return Key.Level>=World->FullResClipLevels&&!GpuFarEligibleTiles.contains(Key);}));
+                    {return Key.Level<World->FullResClipLevels;}));
-            std::printf("[GpuFarEligibility] eligible_gpu_tiles=%zu cpu_transition_fallback_tiles=%zu\n",
+            std::printf("[GpuFieldEligibility] eligible_gpu_tiles=%zu near_gpu_tiles=%zu cpu_fallback_tiles=%zu\n",
-                GpuFarEligibleTiles.size(),CpuTransitionFallbacks);
+                GpuFarEligibleTiles.size(),NearGpuTiles,Keys.size()-GpuFarEligibleTiles.size());
         }
 
         void PumpGpuFarGeneration()
         {
             if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld()||!Device||!Context
                 ||GpuFarQueue.empty())return;
+            const auto PumpStart=std::chrono::steady_clock::now();
+            const double PumpMs=std::chrono::duration<double,std::milli>(PumpStart-ColdProcessStart).count();
+            if(GpuFarFirstPumpMs<0.0)GpuFarFirstPumpMs=PumpMs;
+            if(LastGpuFarPumpAt.time_since_epoch().count()!=0)
+            {
+                const double GapMs=std::chrono::duration<double,std::milli>(PumpStart-LastGpuFarPumpAt).count();
+                GpuFarPumpGapMs+=GapMs;GpuFarPumpGapMaxMs=std::max(GpuFarPumpGapMaxMs,GapMs);
+            }
+            LastGpuFarPumpAt=PumpStart;GpuFarLastPumpMs=PumpMs;++GpuFarPumpCalls;
             const bool bDeferredBatch=!bGpuFarAuditEnabled;
             if(bDeferredBatch)
             {
@@ -2413,8 +2510,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     AuditTile.Indices=std::move(AuditMesh.indices);
                     GpuFarAudit->AuditTile(AuditTile);
                 }
-                const uint64 Cells=static_cast<uint64>(World->CoarseTileCells);
+                Tile.VertexCapacityBytes=static_cast<uint64>(
-                Tile.VertexCapacityBytes=Cells*Cells*Cells*15u*32u;
+                    VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile)*32u;
                 Tile.GpuDispatchMs=GpuDispatchMs;
                 GpuFarEnqueueWallSeconds+=EnqueueMs/1000.0;
                 if(GpuDispatchMs>0.0)
@@ -2451,6 +2548,11 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                     }
                     else
                     {
+                        Batch.SubmittedAt=std::chrono::steady_clock::now();
+                        const double SubmitMs=std::chrono::duration<double,std::milli>(
+                            Batch.SubmittedAt-ColdProcessStart).count();
+                        if(GpuFarFirstSubmitMs<0.0)GpuFarFirstSubmitMs=SubmitMs;
+                        GpuFarLastSubmitMs=SubmitMs;
                         const uint32 BatchIndex=static_cast<uint32>(GpuFarBatches.size());
                         GpuFarBatches.push_back(Batch);
                         for(const FVoxelTileKey& Key:BuiltKeys)
@@ -2471,6 +2573,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 }
             }
             SampleGpuFarMemory();
+            const double PumpWallMs=std::chrono::duration<double,std::milli>(
+                std::chrono::steady_clock::now()-PumpStart).count();
+            GpuFarPumpWallMs+=PumpWallMs;GpuFarPumpMaxWallMs=std::max(GpuFarPumpMaxWallMs,PumpWallMs);
         }
 
         void ToggleFarRenderer()
@@ -2505,7 +2610,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             FCrossingWork& Work=CrossingWork[CrossingIndex];
             Work.SubmittedJobs=Update.NewJobs;Work.PendingJobs=Update.NewJobs;
             Work.StartTime=std::chrono::steady_clock::now();
-            std::printf("[FarRenderer] mode=%s cpu_far_jobs=%d gpu_far_ready=%zu gpu_far_queued=%zu (H toggles)\n",
+            std::printf("[FarRenderer] mode=%s cpu_field_jobs=%d gpu_field_ready=%zu gpu_field_queued=%zu (H toggles)\n",
                 bGpuFarEnabled?"GPU":"CPU",Update.NewJobs,GpuFarTiles.size(),GpuFarQueue.size());
         }
 
@@ -2541,8 +2646,6 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
 
             const double FullShapeSeconds=std::chrono::duration<double>(
                 std::chrono::steady_clock::now()-FullShapeBuildStart).count();
-            std::vector<FVoxelTileKey> ResidentKeys;ResidentKeys.reserve(ResidentTiles.size());
-            for(const auto& Pair:ResidentTiles)ResidentKeys.push_back(Pair.first);
             if(Streamer)Streamer->Stop();
             Streamer.reset();
             ShapeWorld=std::move(PendingFullShapeWorld);
@@ -2551,6 +2654,24 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             const FVector CameraVoxels(Position.X*4.0,Position.Y*4.0,Position.Z*4.0);
             Streamer=std::make_unique<FStandaloneStreamingGenerator>(*World,WorkerCount,false,nullptr,-1,
                 false,false,false,false,ShapeWorld.get());
+            if(bGpuFarEnabled&&ShapeWorld->hasFarWorld())
+            {
+                ScheduleGpuFarTiles(DesiredTiles);
+                if(bGpuFarAuditEnabled)
+                {
+                    std::vector<FVoxelTileKey> AuditKeys;AuditKeys.reserve(GpuFarEligibleTiles.size());
+                    for(const FVoxelTileKey& Key:DesiredTiles)
+                        if(Key.Level>=6&&GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key);
+                    GpuFarAudit=std::make_unique<FFarSurvivalAudit>(*World,*ShapeWorld,AuditKeys,false);
+                }
+                std::vector<FVoxelTileKey> CpuResidentNear;
+                for(const auto& Pair:ResidentTiles)
+                    if(Pair.first.Level<World->FullResClipLevels
+                        &&GpuFarEligibleTiles.contains(Pair.first))CpuResidentNear.push_back(Pair.first);
+                for(const FVoxelTileKey& Key:CpuResidentNear)RemoveResidentTile(Key);
+            }
+            std::vector<FVoxelTileKey> ResidentKeys;ResidentKeys.reserve(ResidentTiles.size());
+            for(const auto& Pair:ResidentTiles)ResidentKeys.push_back(Pair.first);
             ++CrossingIndex;
             const FStandaloneStreamingUpdate Update=Streamer->Start(CameraVoxels,ResidentKeys,CrossingIndex);
             if(Update.DesiredTiles!=static_cast<int32>(DesiredTiles.size()))
@@ -2847,13 +2968,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             }
         }
 
-        bool IsFullDetailUploaded()const
+        bool IsFullDetailUploaded()
         {
             if(!Streamer||World->FullResClipLevels<=0)return false;
-            int32 ResidentFullRes=0;
+            int32 DesiredFullRes=0,ReadyFullRes=0;
             for(const FVoxelTileKey& Key:DesiredTiles)
-                if(Key.Level<World->FullResClipLevels&&ResidentTiles.contains(Key))++ResidentFullRes;
+            {
-            if(ResidentFullRes!=Streamer->FullResDesiredCount())return false;
+                if(Key.Level>=World->FullResClipLevels)continue;
+                ++DesiredFullRes;
+                if(IsRenderedTileReady(Key))++ReadyFullRes;
+            }
+            if(ReadyFullRes!=DesiredFullRes)return false;
             for(const FStandaloneBatchKey& Key:DirtyBatchSet)if(Key.Level<World->FullResClipLevels)return false;
             return true;
         }
@@ -3041,8 +3166,13 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             {
                 bFullDetailReported=true;
                 const double Elapsed=std::chrono::duration<double>(std::chrono::steady_clock::now()-ColdProcessStart).count();
-                std::printf("[FullDetailReady] %.3f s after cold process start; resident=%zu/%d tiles through LOD%d.\n",
+                const std::size_t DesiredNear=static_cast<std::size_t>(std::count_if(
-                    Elapsed,ResidentTiles.size(),Streamer->FullResDesiredCount(),World->FullResClipLevels-1);
+                    DesiredTiles.begin(),DesiredTiles.end(),[&](const FVoxelTileKey& Key)
+                        {return Key.Level<World->FullResClipLevels;}));
+                std::printf("[FullDetailReady] %.3f s after cold process start; rendered=%zu/%zu tiles through LOD%d.\n",
+                    Elapsed,DesiredNear,DesiredNear,World->FullResClipLevels-1);
+                std::printf("[ColdTimeline] event=near_all_tiles_ready at_ms=%.3f waits_for=all_%zu_LOD0_1_gpu_or_cpu_tiles\n",
+                    Elapsed*1000.0,DesiredNear);
                 ReportTileTimings();std::fflush(stdout);
                 StartFarShapeGeneration();
             }
@@ -3056,6 +3186,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 std::printf("[InitialClipmapReady] %.3f s after cold process start; %zu tiles, %llu vertices, %llu indices, %zu GPU batches.\n",
                     Elapsed,DesiredResidentCount,static_cast<unsigned long long>(TotalVertices),
                     static_cast<unsigned long long>(TotalIndices),StreamingBatches.size());
+                std::printf("[ColdTimeline] event=whole_clipmap_ready at_ms=%.3f waits_for=%zu_GPU_tiles_plus_CPU_fallback_and_uploads\n",
+                    Elapsed*1000.0,GpuFarEligibleTiles.size());
                 if(bGpuFarEnabled)
                 {
                     SampleGpuFarMemory();
@@ -3068,6 +3200,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                         GpuFarEnqueueWallSeconds*1000.0,
                         GpuFarGpuDispatchMs,AverageGpuMs,GpuFarGpuDispatchMaxMs,
                         GpuFarSubmissionSpanMs,GpuFarSubmissionSpanMaxMs);
+                    std::printf("[GpuFarTimeline] pump_calls=%llu pump_first_ms=%.3f pump_last_ms=%.3f pump_gap_sum_ms=%.3f pump_gap_max_ms=%.3f pump_cpu_sum_ms=%.3f pump_cpu_max_ms=%.3f submit_first_ms=%.3f submit_last_ms=%.3f completion_polls=%llu completion_first_ms=%.3f completion_last_ms=%.3f queue_envelope_ms=%.3f gpu_busy_ms=%.3f estimated_idle_or_poll_ms=%.3f\n",
+                        static_cast<unsigned long long>(GpuFarPumpCalls),GpuFarFirstPumpMs,GpuFarLastPumpMs,
+                        GpuFarPumpGapMs,GpuFarPumpGapMaxMs,GpuFarPumpWallMs,GpuFarPumpMaxWallMs,
+                        GpuFarFirstSubmitMs,GpuFarLastSubmitMs,
+                        static_cast<unsigned long long>(GpuFarCompletionPolls),GpuFarFirstCompletionPollMs,
+                        GpuFarLastCompletionPollMs,
+                        GpuFarFirstSubmitMs>=0.0&&GpuFarLastCompletionPollMs>=0.0
+                            ?GpuFarLastCompletionPollMs-GpuFarFirstSubmitMs:0.0,
+                        GpuFarSubmissionSpanMs,
+                        GpuFarFirstSubmitMs>=0.0&&GpuFarLastCompletionPollMs>=0.0
+                            ?std::max(0.0,GpuFarLastCompletionPollMs-GpuFarFirstSubmitMs-GpuFarSubmissionSpanMs):0.0);
                     std::printf("[GpuFarMemory] dxgi_local_sampled_peak_mb=%.1f dxgi_local_budget_mb=%.1f\n",
                         GpuFarLocalMemoryPeakBytes/(1024.0*1024.0),
                         GpuFarLocalMemoryBudgetBytes/(1024.0*1024.0));
@@ -3114,7 +3257,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             uint64 VisibleIndices=0;for(const auto& Pair:StreamingBatches)if(bShowLevel[Pair.first.Level])VisibleIndices+=Pair.second.IndexCount;
             std::ostringstream S;S<<"VoxelForge standalone | seed "<<Seed<<" | "<<(bShapeWorld?"shape-interval owner world":(bRoundTunnels?"round/winding research tunnels":"base world"))
                 <<(bRoundTunnels?(bRoundTunnelEnhanced?" (detailed walls)":" (legacy walls)"):"")
-                <<" | "<<VisibleIndices/3<<" tris | "<<FrameMs<<" ms | far "<<(bGpuFarEnabled?"GPU":"CPU")
+                <<" | "<<VisibleIndices/3<<" tris | "<<FrameMs<<" ms | field "<<(bGpuFarEnabled?"GPU":"CPU")
                 <<" | "<<(bGenerationRunning?"generating":"ready")
                 <<" | pos "<<std::fixed<<std::setprecision(1)<<Position.X<<", "<<Position.Y<<", "<<Position.Z<<" m";
             SetWindowTextA(Window,S.str().c_str());
@@ -3122,8 +3265,13 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
 
         void Draw()
         {
+            const auto DrawStart=std::chrono::steady_clock::now();
+            const double InterDrawGapMs=LastDrawEndAt.time_since_epoch().count()!=0
+                ?std::chrono::duration<double,std::milli>(DrawStart-LastDrawEndAt).count():0.0;
             PumpGeneration();
+            const auto CpuPumpEnd=std::chrono::steady_clock::now();
             PumpGpuFarGeneration();
+            const auto GpuPumpEnd=std::chrono::steady_clock::now();
             const auto Now=std::chrono::steady_clock::now();
             if(FlightTargetCrossings>0&&bFullDetailReported&&!bFlightStarted)
             {
@@ -3179,7 +3327,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
             // Avoid issuing hundreds of tiny indirect draw calls while the
             // initial far batches are still being built. Once the clipmap is
             // ready the pooled GPU meshes draw directly as before.
-            if(bGpuFarEnabled&&bInitialClipmapReported&&ShapeWorld&&ShapeWorld->hasFarWorld())
+            if(bGpuFarEnabled&&ShapeWorld&&ShapeWorld->hasFarWorld())
             {
                 C.Padding[0]=0.01f;
                 Context->UpdateSubresource(Constants,0,nullptr,&C,0,0);
@@ -3188,6 +3336,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 {
                     const FGpuFarTile& Tile=Pair.second;
                     if(!bShowLevel[Pair.first.Level]||!Tile.VertexBuffer||!Tile.IndexBuffer||!Tile.IndirectArgs)continue;
+                    if(!bInitialClipmapReported&&Pair.first.Level>=World->FullResClipLevels)continue;
+                    if(!IsRenderedTileReady(Pair.first))continue;
                     const UINT Stride=32,Offset=0;
                     Context->IASetVertexBuffers(0,1,&Tile.VertexBuffer,&Stride,&Offset);
                     Context->IASetIndexBuffer(Tile.IndexBuffer,DXGI_FORMAT_R32_UINT,0);
@@ -3202,7 +3352,25 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                 }
             }
             CaptureRenderedFrame();
-            SwapChain->Present(1,0);UpdateTitle();
+            const auto BeforePresent=std::chrono::steady_clock::now();
+            const UINT PresentSyncInterval=1u;
+            const auto PresentCallStart=std::chrono::steady_clock::now();
+            SwapChain->Present(PresentSyncInterval,0);
+            const auto PresentCallEnd=std::chrono::steady_clock::now();
+            UpdateTitle();
+            const auto DrawEnd=std::chrono::steady_clock::now();
+            const double CpuPumpMs=std::chrono::duration<double,std::milli>(CpuPumpEnd-DrawStart).count();
+            const double GpuPumpMs=std::chrono::duration<double,std::milli>(GpuPumpEnd-CpuPumpEnd).count();
+            const double PresentCallMs=std::chrono::duration<double,std::milli>(PresentCallEnd-PresentCallStart).count();
+            const double TitleMs=std::chrono::duration<double,std::milli>(DrawEnd-PresentCallEnd).count();
+            const double RenderPresentMs=PresentCallMs+TitleMs;
+            if(InterDrawGapMs>50.0||CpuPumpMs>50.0||GpuPumpMs>50.0||RenderPresentMs>50.0)
+                std::printf("[LongViewerFrame] at_ms=%.3f inter_draw_gap_ms=%.3f cpu_pump_ms=%.3f gpu_pump_ms=%.3f pre_present_ms=%.3f present_sync=%u present_call_ms=%.3f title_ms=%.3f\n",
+                    std::chrono::duration<double,std::milli>(DrawStart-ColdProcessStart).count(),
+                    InterDrawGapMs,CpuPumpMs,GpuPumpMs,
+                    std::chrono::duration<double,std::milli>(BeforePresent-GpuPumpEnd).count(),
+                    PresentSyncInterval,PresentCallMs,TitleMs);
+            LastDrawEndAt=DrawEnd;
         }
     };
 
@@ -5106,6 +5274,8 @@ int main(int Argc,char** Argv)
     }
     App.bProfileDetailBand=bProfileDetailBand;
     std::string Error;
+    std::printf("[ColdTimeline] stage=owner_data_load begin_ms=%.3f\n",
+        std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count());
     App.World=std::make_shared<FStandaloneOwnerWorld>();
     {
         std::ifstream JsonFile(App.DataPath,std::ios::binary);
@@ -5118,6 +5288,8 @@ int main(int Argc,char** Argv)
     }
     if(!StandaloneOwnerLoader::Resolve(App.OwnerJsonRoot,*App.World,Error))
     {std::fprintf(stderr,"Owner data load failed: %s\n",Error.c_str());return 2;}
+    std::printf("[ColdTimeline] stage=owner_data_load end_ms=%.3f\n",
+        std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count());
     App.Seed=App.World->Seed;
     if(App.bParityOnly)
     {
@@ -5213,7 +5385,13 @@ int main(int Argc,char** Argv)
     App.Window=CreateWindowA(Class.lpszClassName,"VoxelForge standalone world",WS_OVERLAPPEDWINDOW,CW_USEDEFAULT,CW_USEDEFAULT,
         Rect.right-Rect.left,Rect.bottom-Rect.top,nullptr,nullptr,Class.hInstance,nullptr);
     GApp=&App;
+    const auto DeviceInitStart=std::chrono::steady_clock::now();
+    std::printf("[ColdTimeline] stage=viewer_device_init begin_ms=%.3f\n",
+        std::chrono::duration<double,std::milli>(DeviceInitStart-ColdProcessStart).count());
     if(!App.Window||!App.InitializeDevice()){MessageBoxA(nullptr,"Direct3D 11 initialization failed.","VoxelForge standalone",MB_ICONERROR);return 3;}
+    std::printf("[ColdTimeline] stage=viewer_device_init end_ms=%.3f duration_ms=%.3f\n",
+        std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-ColdProcessStart).count(),
+        std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-DeviceInitStart).count());
     ShowWindow(App.Window,SW_SHOW);UpdateWindow(App.Window);
     std::printf("Controls: hold right mouse to look; WASD move; Space/Ctrl (E/Q) up/down; Shift faster; wheel changes speed; F wireframe; G toggles legacy/detailed round walls; H toggles CPU/GPU far tiles; 1-9 toggle LOD; 0 toggle all levels; N/P change seed; T swaps base/round tunnels; V swaps owner density/shape intervals; Esc quit.\n");
     std::printf("Owner data: %s; starting at first-strate open-space sample (%.2f, %.2f, %.2f) m.\n",