wip(standalone): GPU step 8 partial -- GPU skirts, first-present hold; remote gap open

Luna's session ended on an API "Bad Request" before a final report; state is unverified by me. In the log: GPU clipped-skirt emitter (count/scan/emit, protected signs and critical air honoured), exposed-skirt audit passing at the end; first-present stall reproduced on every launch (2.36-2.59 s, not exe identity/idle/CPU run) and hidden by holding the CPU-ready frame until GPU residency. Remote LOD3-4 band still differs from CPU; draw ranges and the one inspected LOD3 tile match, cause not yet found (one-tile probe mid-fix). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01MmSoF64BbRFmk9C3p1VsBJ
2026-10-07 02:19:38 +02:00
1 parent 90421d0655
commit f2de8fe9c6
@@ -255,6 +255,11 @@ namespace VoxelForgeCore
        Vec3i origin{};
        std::int32_t cells = 0, step = 0, sampleDim = 0;
        bool nearDetailEnabled = false;
+        // CPU-decided clipmap seam data. The GPU evaluates the ordinary far-field
+        // rock/air signs used to clip skirts; this mask only says which tile faces
+        // need a curtain, and the depth mirrors the CPU mesher in voxel units.
+        std::uint8_t visibleSkirtFaceMask = 0;
+        float skirtDepthVoxels = 0.0f;
        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
@@ -120,7 +120,7 @@ if ($Target -eq 'All') {
        Sort-Object -Property FullName -Descending | Select-Object -First 1
    if ($FxC) {
        [IO.Directory]::CreateDirectory($PrecompiledShaderDir) | Out-Null
-        foreach ($Entry in @('EvaluateFarFieldCS','ClassifyCS','ClassifyFarCS','ScanGroupsCS','EmitCS','NoiseProbeCS','BuildFarConnectorRingsCS','EmitFarConnectorPanelsCS')) {
+        foreach ($Entry in @('EvaluateFarFieldCS','ClassifyCS','ClassifyFarCS','ScanGroupsCS','EmitCS','NoiseProbeCS','BuildFarConnectorRingsCS','EmitFarConnectorPanelsCS','ClassifyFarSkirtsCS','ScanFarSkirtsCS','EmitFarSkirtsCS')) {
            $ShaderBinary = Join-Path $PrecompiledShaderDir ($Entry+'.cso')
            $CommandLine = '"{0}" /nologo /T cs_5_0 /E {1} /O3 /Gis /Ges /Fo "{2}" "{3}"' -f $FxC.FullName,$Entry,$ShaderBinary,$ShaderSource
            Invoke-FieldViewerBuildCommand $CommandLine "precompile HLSL entry $Entry"
@@ -0,0 +1,70 @@
+import argparse
+from pathlib import Path
+
+import numpy as np
+from PIL import Image
+
+
+ROOT = Path(__file__).resolve().parent
+parser = argparse.ArgumentParser()
+parser.add_argument("--prefix", default="step8")
+args = parser.parse_args()
+if not args.prefix or any(c not in "abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789_-" for c in args.prefix):
+    raise SystemExit("prefix must contain only letters, digits, underscores, or hyphens")
+VIEWS = (
+    "chamber",
+    "tunnel",
+    "passage",
+    "far",
+    "passage_floor_closeup",
+    "tunnel_bend",
+    "remote_yaw0",
+    "remote_yaw90",
+    "remote_yaw180",
+    "remote_yaw270",
+)
+
+rows = []
+for view in VIEWS:
+    cpu_path = ROOT / f"{args.prefix}_cpu_{view}.png"
+    gpu_path = ROOT / f"{args.prefix}_gpu_{view}.png"
+    cpu = np.asarray(Image.open(cpu_path).convert("RGB"), dtype=np.int16)
+    gpu = np.asarray(Image.open(gpu_path).convert("RGB"), dtype=np.int16)
+    if cpu.shape != gpu.shape:
+        raise RuntimeError(f"dimension mismatch for {view}: {cpu.shape} vs {gpu.shape}")
+    diff = np.abs(cpu - gpu)
+    pixel_max = diff.max(axis=2)
+    rows.append(
+        {
+            "view": view,
+            "size": f"{cpu.shape[1]}x{cpu.shape[0]}",
+            "mae_rgb": float(diff.mean()),
+            "rmse_rgb": float(np.sqrt(np.square(diff, dtype=np.float64).mean())),
+            "p95_pixel_max": float(np.percentile(pixel_max, 95)),
+            "p99_pixel_max": float(np.percentile(pixel_max, 99)),
+            "max_pixel_channel": int(pixel_max.max()),
+            "changed_pct": float(np.count_nonzero(pixel_max) * 100.0 / pixel_max.size),
+            "over_8_pct": float(np.count_nonzero(pixel_max > 8) * 100.0 / pixel_max.size),
+            "over_32_pct": float(np.count_nonzero(pixel_max > 32) * 100.0 / pixel_max.size),
+        }
+    )
+
+lines = [
+    f"# {args.prefix} 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 |",
+    "|---|---:|---:|---:|---:|---:|---:|---:|---:|",
+]
+for row in rows:
+    lines.append(
+        f"| {row['view']} | {row['mae_rgb']:.4f} | {row['rmse_rgb']:.4f} | "
+        f"{row['p95_pixel_max']:.0f} | {row['p99_pixel_max']:.0f} | {row['max_pixel_channel']} | "
+        f"{row['changed_pct']:.2f}% | {row['over_8_pct']:.2f}% | {row['over_32_pct']:.2f}% |"
+    )
+
+report = ROOT / f"{args.prefix}_capture_comparison.md"
+report.write_text("\n".join(lines) + "\n", encoding="utf-8")
+print("\n".join(lines))
+print(f"\nWrote {report}")
@@ -0,0 +1,50 @@
+from pathlib import Path
+
+import numpy as np
+from PIL import Image
+
+
+ROOT = Path(__file__).resolve().parent
+rows = []
+for lod in range(2, 9):
+    cpu = np.asarray(
+        Image.open(ROOT / f"step8_lod{lod}_cpu_remote_yaw0.png").convert("RGB"),
+        dtype=np.int16,
+    )
+    gpu = np.asarray(
+        Image.open(ROOT / f"step8_lod{lod}_gpu_remote_yaw0.png").convert("RGB"),
+        dtype=np.int16,
+    )
+    if cpu.shape != gpu.shape:
+        raise RuntimeError(f"LOD {lod} dimensions differ: {cpu.shape} vs {gpu.shape}")
+    diff = np.abs(cpu - gpu)
+    pixel_max = diff.max(axis=2)
+    rows.append(
+        {
+            "lod": lod,
+            "mae": float(diff.mean()),
+            "rmse": float(np.sqrt(np.square(diff, dtype=np.float64).mean())),
+            "changed": float(np.count_nonzero(pixel_max) * 100.0 / pixel_max.size),
+            "over32": float(np.count_nonzero(pixel_max > 32) * 100.0 / pixel_max.size),
+            "max": int(pixel_max.max()),
+        }
+    )
+
+lines = [
+    "# Step 8 remote yaw 0, isolated LOD render comparison",
+    "",
+    "The viewer generated and fully-resident clip radius 3 for each pair; `--capture-only-lod` filtered drawing only.",
+    "",
+    "| LOD | MAE RGB | RMSE RGB | Changed pixels | Pixels >32 | Max channel |",
+    "|---:|---:|---:|---:|---:|---:|",
+]
+for row in rows:
+    lines.append(
+        f"| {row['lod']} | {row['mae']:.4f} | {row['rmse']:.4f} | "
+        f"{row['changed']:.2f}% | {row['over32']:.2f}% | {row['max']} |"
+    )
+
+report = ROOT / "step8_lod_isolation_comparison.md"
+report.write_text("\n".join(lines) + "\n", encoding="utf-8")
+print("\n".join(lines))
+print(f"\nWrote {report}")
@@ -26,6 +26,7 @@
 #include <stdexcept>
 #include <thread>
 #include <unordered_map>
+#include <unordered_set>
 #include <utility>
 #include <vector>

@@ -111,7 +112,9 @@ namespace
    struct FarConnectorConstants
    {
        std::uint32_t connectorCount=0,ringCount=0,vertexCount=0,panelVertexCapacity=0;
-        std::uint32_t shapeCount=0,padding[3]{};
+        std::uint32_t shapeCount=0,skirtFaceMask=0;
+        float skirtDepthVoxels=0.0f;
+        std::uint32_t padding=0;
    };
    static_assert(sizeof(FarConnectorConstants)==32);

@@ -120,6 +123,8 @@ namespace
        float position[3];
        float normal[3];
        float uv[2];
+        friend bool operator==(const GpuVertex& a,const GpuVertex& b) noexcept
+        {return std::memcmp(&a,&b,sizeof(GpuVertex))==0;}
    };
    static_assert(sizeof(GpuVertex) == 32);

@@ -179,12 +184,14 @@ namespace
    struct GpuResult
    {
        std::vector<GpuVertex> vertices;
+        std::vector<GpuVertex> skirtVertices;
        std::vector<std::uint32_t> cases;
        std::vector<float> fieldDensity;
        ComPtr<ID3D11Buffer> vertexBuffer, indexBuffer, indirectArgs;
        ComPtr<ID3D11Buffer> connectorVertexBuffer,connectorIndexBuffer;
        ComPtr<ID3D11Query> startTimestamp,endTimestamp,disjointTimestamp;
        std::array<std::uint32_t,5> drawArgs{};
+        std::array<std::uint32_t,6> connectorDrawArgs{};
        bool sequentialIndices = false;
        bool indexedIndirectArgsValid = false;
        std::uint32_t outputBaseVertex = 0, outputArgsByteOffset = 0;
@@ -195,7 +202,7 @@ namespace
        std::uint64_t d3dBufferBytes = 0;
        double dispatchMs = 0.0;
        double fieldMs = 0.0;
-        std::array<double,6> kernelMs{};
+        std::array<double,9> kernelMs{};
        std::uint64_t fieldDistanceEvaluations=0;
        std::uint32_t maxFieldDistanceEvaluationsPerSample=0;
        double wallMs = 0.0;
@@ -506,6 +513,7 @@ namespace
        struct FarOutputChunk
        {
            ComPtr<ID3D11Buffer> vertices, indices, args;
+            ComPtr<ID3D11ShaderResourceView> surfaceVertexSrv;
            ComPtr<ID3D11UnorderedAccessView> vertexUav, indexUav, argsUav;
            ComPtr<ID3D11Buffer> connectorVertices,connectorIndices;
            ComPtr<ID3D11UnorderedAccessView> connectorVertexUav,connectorIndexUav;
@@ -528,6 +536,9 @@ namespace
            ComPtr<ID3D11ShaderResourceView> connectorRingSrv;
            ComPtr<ID3D11UnorderedAccessView> densityUav,nearNormalDensityUav,caseUav,localOffsetUav,groupCountUav,groupOffsetUav;
            ComPtr<ID3D11UnorderedAccessView> connectorRingUav;
+            ComPtr<ID3D11Buffer> skirtScratch;
+            ComPtr<ID3D11UnorderedAccessView> skirtScratchUav;
+            std::uint32_t skirtScratchCapacity=0;
            DynamicInput shapes,ranges,ids,protectedAir,protectedSolid,nearPackedProxy;
            DynamicInput connectors;
            std::uint32_t connectorRingCapacity=0;
@@ -542,6 +553,7 @@ namespace
            ComPtr<ID3D11DeviceContext> immediateContext,deferredContext;
        ComPtr<ID3D11ComputeShader> fieldEval, classify, fieldClassify, scan, emit, noiseProbe;
        ComPtr<ID3D11ComputeShader> connectorRings,connectorPanels;
+        ComPtr<ID3D11ComputeShader> classifyFarSkirts,scanFarSkirts,emitFarSkirts;
        ComPtr<IDXGIAdapter3> memoryAdapter;
        std::string name;
        UINT64 timestampFrequency = 0;
@@ -550,7 +562,10 @@ namespace
        std::uint32_t farOutputSlotsPerChunk=128,farMaxVerticesPerTile=0;
        static constexpr std::uint32_t FarOutputVertexCapacityPerTile=
            VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile;
-        static constexpr std::uint32_t FarPanelVertexCapacity=28160;
+        // Full clip-radius-3 skirt audit peaks at 27,840 panel+skirt vertices per
+        // far tile; this leaves 4,928 records of headroom without allocating 49k
+        // records for every near and far output slot.
+        static constexpr std::uint32_t FarPanelVertexCapacity=32768;

        void CompileShaders(const std::filesystem::path& shaderPath)
        {
@@ -586,6 +601,9 @@ namespace
            Create("NoiseProbeCS",noiseProbe);
            Create("BuildFarConnectorRingsCS",connectorRings);
            Create("EmitFarConnectorPanelsCS",connectorPanels);
+            Create("ClassifyFarSkirtsCS",classifyFarSkirts);
+            Create("ScanFarSkirtsCS",scanFarSkirts);
+            Create("EmitFarSkirtsCS",emitFarSkirts);
            std::printf("[GpuPipelineProfile] stage=shader_pipeline_total duration_ms=%.3f\n",
                std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-PipelineStart).count());
        }
@@ -783,17 +801,18 @@ namespace
                    farOutputCapacitySlots-chunkIndex*farOutputSlotsPerChunk);
                const std::uint64_t vertexBytes=static_cast<std::uint64_t>(slots)*farMaxVerticesPerTile*sizeof(GpuVertex);
                const std::uint64_t indexBytes=static_cast<std::uint64_t>(slots)*farMaxVerticesPerTile*sizeof(std::uint32_t);
-                const std::uint64_t argsBytes=static_cast<std::uint64_t>(slots)*10u*sizeof(std::uint32_t);
+                const std::uint64_t argsBytes=static_cast<std::uint64_t>(slots)*11u*sizeof(std::uint32_t);
                if(vertexBytes>std::numeric_limits<UINT>::max()||indexBytes>std::numeric_limits<UINT>::max())
                    throw std::runtime_error("far output chunk exceeds the D3D11 buffer size limit");
                FarOutputChunk chunk;
                chunk.vertices=RawMeshBuffer(static_cast<UINT>(vertexBytes),
-                    D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_VERTEX_BUFFER);
+                    D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_VERTEX_BUFFER|D3D11_BIND_SHADER_RESOURCE);
                chunk.indices=RawMeshBuffer(static_cast<UINT>(indexBytes),
                    D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_INDEX_BUFFER);
                chunk.args=RawMeshBuffer(static_cast<UINT>(argsBytes),D3D11_BIND_UNORDERED_ACCESS,
                    D3D11_RESOURCE_MISC_BUFFER_ALLOW_RAW_VIEWS|D3D11_RESOURCE_MISC_DRAWINDIRECT_ARGS);
                chunk.vertexUav=RawUav(chunk.vertices.Get(),static_cast<UINT>(vertexBytes/sizeof(std::uint32_t)));
+                chunk.surfaceVertexSrv=RawSrv(chunk.vertices.Get(),static_cast<UINT>(vertexBytes/sizeof(std::uint32_t)));
                chunk.indexUav=RawUav(chunk.indices.Get(),static_cast<UINT>(indexBytes/sizeof(std::uint32_t)));
                chunk.argsUav=RawUav(chunk.args.Get(),static_cast<UINT>(argsBytes/sizeof(std::uint32_t)));
                const std::uint64_t connectorVertexBytes=static_cast<std::uint64_t>(slots)
@@ -885,6 +904,9 @@ namespace
            scratch.localOffsets=Structured<std::uint32_t>(nullptr,cells,D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_SHADER_RESOURCE);
            scratch.groupCounts=Structured<std::uint32_t>(nullptr,groups,D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_SHADER_RESOURCE);
            scratch.groupOffsets=Structured<std::uint32_t>(nullptr,groups,D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_SHADER_RESOURCE);
+            scratch.skirtScratchCapacity=std::max<std::uint32_t>(1,farMaxVerticesPerTile);
+            scratch.skirtScratch=Structured<std::uint32_t>(nullptr,scratch.skirtScratchCapacity,
+                D3D11_BIND_UNORDERED_ACCESS);
            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);
@@ -901,6 +923,7 @@ namespace
            scratch.localOffsetUav=Uav(scratch.localOffsets.Get(),cells);
            scratch.groupCountUav=Uav(scratch.groupCounts.Get(),groups);
            scratch.groupOffsetUav=Uav(scratch.groupOffsets.Get(),groups);
+            scratch.skirtScratchUav=Uav(scratch.skirtScratch.Get(),scratch.skirtScratchCapacity);
            const auto DynamicConstant=[&](UINT bytes,const char* label)
            {
                D3D11_BUFFER_DESC desc{};desc.ByteWidth=bytes;desc.Usage=D3D11_USAGE_DYNAMIC;
@@ -938,6 +961,19 @@ namespace
            return view;
        }

+        ComPtr<ID3D11ShaderResourceView> RawSrv(ID3D11Buffer* buffer,UINT dwordCount)
+        {
+            D3D11_SHADER_RESOURCE_VIEW_DESC desc{};
+            desc.Format=DXGI_FORMAT_R32_TYPELESS;
+            desc.ViewDimension=D3D11_SRV_DIMENSION_BUFFEREX;
+            desc.BufferEx.NumElements=dwordCount;
+            desc.BufferEx.Flags=D3D11_BUFFEREX_SRV_FLAG_RAW;
+            ComPtr<ID3D11ShaderResourceView> view;
+            Check(device->CreateShaderResourceView(buffer,&desc,&view),
+                "CreateShaderResourceView(raw surface vertices)");
+            return view;
+        }
+
        void WaitForQuery(ID3D11Query* query)
        {
            for (;;)
@@ -998,7 +1034,7 @@ namespace
                const std::uint32_t localSlot=outputSlot%farOutputSlotsPerChunk;
                pooledChunk=&farOutputChunks[chunkIndex];
                outputBaseVertex=localSlot*farMaxVerticesPerTile;
-                outputArgsByteOffset=localSlot*10u*sizeof(std::uint32_t);
+                outputArgsByteOffset=localSlot*11u*sizeof(std::uint32_t);
            }
            if (farField && (farField->sampleDim != static_cast<std::int32_t>(tile.dim)
                || farField->cells != static_cast<std::int32_t>(tile.cells)
@@ -1153,36 +1189,49 @@ namespace
            }
            ComPtr<ID3D11Buffer> args,vertices,indices;
            ComPtr<ID3D11UnorderedAccessView> argsUav,vertexUav,indexUav;
+            ComPtr<ID3D11ShaderResourceView> surfaceVertexSrv;
            ComPtr<ID3D11Buffer> connectorVertices,connectorIndices;
            ComPtr<ID3D11UnorderedAccessView> connectorVertexUav,connectorIndexUav;
+            ComPtr<ID3D11Buffer> skirtScratchBuffer;
+            ComPtr<ID3D11UnorderedAccessView> skirtScratchUav;
            const std::uint32_t connectorPanelCapacity=pooledChunk?FarPanelVertexCapacity:
-                std::max<std::uint32_t>(1,farField?farField->connectorVertexCount:0);
+                std::max<std::uint32_t>(1,farField&&(farField->visibleSkirtFaceMask!=0)
+                    ?FarPanelVertexCapacity:(farField?farField->connectorVertexCount:0));
            if(farField&&farField->connectorVertexCount>connectorPanelCapacity)
                throw std::runtime_error("far connector repair output exceeds the pooled panel capacity");
            if(pooledChunk)
            {
                args=pooledChunk->args;vertices=pooledChunk->vertices;indices=pooledChunk->indices;
+                surfaceVertexSrv=pooledChunk->surfaceVertexSrv;
                argsUav=pooledChunk->argsUav;vertexUav=pooledChunk->vertexUav;indexUav=pooledChunk->indexUav;
                connectorVertices=pooledChunk->connectorVertices;connectorIndices=pooledChunk->connectorIndices;
                connectorVertexUav=pooledChunk->connectorVertexUav;connectorIndexUav=pooledChunk->connectorIndexUav;
+                skirtScratchBuffer=farScratch.skirtScratch;skirtScratchUav=farScratch.skirtScratchUav;
            }
            else
            {
-                args = RawMeshBuffer(10 * sizeof(std::uint32_t), D3D11_BIND_UNORDERED_ACCESS,
+                args = RawMeshBuffer(11 * sizeof(std::uint32_t), D3D11_BIND_UNORDERED_ACCESS,
                    D3D11_RESOURCE_MISC_BUFFER_ALLOW_RAW_VIEWS | D3D11_RESOURCE_MISC_DRAWINDIRECT_ARGS);
                vertices = RawMeshBuffer(static_cast<UINT>(maxVertices * sizeof(GpuVertex)),
-                    D3D11_BIND_UNORDERED_ACCESS | D3D11_BIND_VERTEX_BUFFER);
+                    D3D11_BIND_UNORDERED_ACCESS | D3D11_BIND_VERTEX_BUFFER | D3D11_BIND_SHADER_RESOURCE);
                indices = RawMeshBuffer(static_cast<UINT>(maxVertices * sizeof(std::uint32_t)),
                    D3D11_BIND_UNORDERED_ACCESS | D3D11_BIND_INDEX_BUFFER);
-                argsUav=RawUav(args.Get(),10);
+                argsUav=RawUav(args.Get(),11);
                vertexUav=RawUav(vertices.Get(),static_cast<UINT>(maxVertices*sizeof(GpuVertex)/sizeof(std::uint32_t)));
                indexUav=RawUav(indices.Get(),static_cast<UINT>(maxVertices));
+                surfaceVertexSrv=RawSrv(vertices.Get(),static_cast<UINT>(maxVertices*sizeof(GpuVertex)/sizeof(std::uint32_t)));
                connectorVertices=RawMeshBuffer(connectorPanelCapacity*sizeof(GpuVertex),
                    D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_VERTEX_BUFFER);
                connectorIndices=RawMeshBuffer(connectorPanelCapacity*sizeof(std::uint32_t),
                    D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_INDEX_BUFFER);
                connectorVertexUav=RawUav(connectorVertices.Get(),connectorPanelCapacity*sizeof(GpuVertex)/sizeof(std::uint32_t));
                connectorIndexUav=RawUav(connectorIndices.Get(),connectorPanelCapacity);
+                if(farField&&farField->visibleSkirtFaceMask!=0)
+                {
+                    skirtScratchBuffer=Structured<std::uint32_t>(nullptr,
+                        static_cast<std::uint32_t>(maxVertices),D3D11_BIND_UNORDERED_ACCESS);
+                    skirtScratchUav=Uav(skirtScratchBuffer.Get(),static_cast<UINT>(maxVertices));
+                }
            }
            D3D11_BUFFER_DESC cbDesc{};
            cbDesc.ByteWidth = sizeof(TileConstants);
@@ -1213,7 +1262,9 @@ namespace
                const FarConnectorConstants connectorConstants{
                    static_cast<std::uint32_t>(farField->connectorPanels.size()),
                    farField->connectorRingCount,farField->connectorVertexCount,connectorPanelCapacity,
-                    static_cast<std::uint32_t>(farField->worldInput->shapes.size()),{0,0,0}};
+                    static_cast<std::uint32_t>(farField->worldInput->shapes.size()),
+                    static_cast<std::uint32_t>(farField->visibleSkirtFaceMask),
+                    farField->skirtDepthVoxels,0};
                if(pooledChunk)
                {
                    cb=farScratch.tileConstants;farCb=farScratch.farConstants;
@@ -1255,7 +1306,7 @@ namespace
            D3D11_QUERY_DESC disjointDesc{D3D11_QUERY_TIMESTAMP_DISJOINT, 0};
            ComPtr<ID3D11Query> startQuery, endQuery, disjointQuery;
            ComPtr<ID3D11Query> fieldStartQuery,fieldEndQuery;
-            std::array<ComPtr<ID3D11Query>,6> kernelStartQueries,kernelEndQueries;
+            std::array<ComPtr<ID3D11Query>,9> kernelStartQueries,kernelEndQueries;
            if(captureGpuTiming)
            {
                Check(device->CreateQuery(&timestampDesc, &startQuery), "CreateQuery(start)");
@@ -1388,6 +1439,49 @@ namespace
                    context->CSSetUnorderedAccessViews(0,7,nullUavs,nullptr);
                }
            }
+            if(farField&&farField->visibleSkirtFaceMask!=0)
+            {
+                ID3D11Buffer* skirtCbs[3]={cb.Get(),farCb.Get(),farConnectorConstantsBuffer.Get()};
+                context->CSSetConstantBuffers(0,3,skirtCbs);
+                ID3D11ShaderResourceView* skirtSrvs[20]{};
+                skirtSrvs[10]=farShapeSrv.Get();skirtSrvs[11]=farRangeSrv.Get();
+                skirtSrvs[12]=farIdSrv.Get();skirtSrvs[13]=farProtectedSrv.Get();
+                skirtSrvs[14]=farProtectedSolidSrv.Get();skirtSrvs[19]=surfaceVertexSrv.Get();
+                context->CSSetShaderResources(0,20,skirtSrvs);
+                ID3D11UnorderedAccessView* classifySkirtUavs[8]{};
+                classifySkirtUavs[0]=skirtScratchUav.Get();classifySkirtUavs[4]=argsUav.Get();
+                const std::uint32_t skirtDispatchVertices=constants.outputVertexCapacity;
+                beginKernel(6);
+                context->CSSetShader(classifyFarSkirts.Get(),nullptr,0);
+                context->CSSetUnorderedAccessViews(0,8,classifySkirtUavs,nullptr);
+                context->Dispatch((skirtDispatchVertices+63u)/64u,1,1);
+                endKernel(6);
+                ID3D11ShaderResourceView* nullSkirtSrvs[20]{};
+                ID3D11UnorderedAccessView* nullSkirtUavs[8]{};
+                context->CSSetShaderResources(0,20,nullSkirtSrvs);
+                context->CSSetUnorderedAccessViews(0,8,nullSkirtUavs,nullptr);
+
+                beginKernel(7);
+                context->CSSetShader(scanFarSkirts.Get(),nullptr,0);
+                ID3D11UnorderedAccessView* scanSkirtUavs[8]{};
+                scanSkirtUavs[0]=skirtScratchUav.Get();scanSkirtUavs[4]=argsUav.Get();
+                context->CSSetUnorderedAccessViews(0,8,scanSkirtUavs,nullptr);
+                context->Dispatch(1,1,1);
+                endKernel(7);
+                context->CSSetUnorderedAccessViews(0,8,nullSkirtUavs,nullptr);
+
+                context->CSSetShaderResources(0,20,skirtSrvs);
+                beginKernel(8);
+                context->CSSetShader(emitFarSkirts.Get(),nullptr,0);
+                ID3D11UnorderedAccessView* emitSkirtUavs[8]{};
+                emitSkirtUavs[0]=skirtScratchUav.Get();emitSkirtUavs[4]=argsUav.Get();
+                emitSkirtUavs[5]=connectorVertexUav.Get();emitSkirtUavs[6]=connectorIndexUav.Get();
+                context->CSSetUnorderedAccessViews(0,8,emitSkirtUavs,nullptr);
+                context->Dispatch((skirtDispatchVertices+63u)/64u,1,1);
+                endKernel(8);
+                context->CSSetShaderResources(0,20,nullSkirtSrvs);
+                context->CSSetUnorderedAccessViews(0,8,nullSkirtUavs,nullptr);
+            }
            if(captureGpuTiming)
            {
                context->End(endQuery.Get());
@@ -1405,11 +1499,13 @@ namespace
                    * sizeof(GpuVertex)));
                auto indexStaging = Staging(static_cast<UINT>(readbackCapacity
                    * sizeof(std::uint32_t)));
-                auto argsStaging = Staging(10 * sizeof(std::uint32_t));
-                auto connectorVertexStaging=farField&&farField->connectorVertexCount
-                    ?Staging(farField->connectorVertexCount*sizeof(GpuVertex)):ComPtr<ID3D11Buffer>{};
-                auto connectorIndexStaging=farField&&farField->connectorVertexCount
-                    ?Staging(farField->connectorVertexCount*sizeof(std::uint32_t)):ComPtr<ID3D11Buffer>{};
+                auto argsStaging = Staging(11 * sizeof(std::uint32_t));
+                const bool hasConnectorOutput=farField&&(farField->connectorVertexCount
+                    ||farField->visibleSkirtFaceMask!=0);
+                auto connectorVertexStaging=hasConnectorOutput
+                    ?Staging(connectorPanelCapacity*sizeof(GpuVertex)):ComPtr<ID3D11Buffer>{};
+                auto connectorIndexStaging=hasConnectorOutput
+                    ?Staging(connectorPanelCapacity*sizeof(std::uint32_t)):ComPtr<ID3D11Buffer>{};
                if (QueryLocalMemory(localUsage, localBudget))
                {
                    result.hasAdapterMemoryInfo = true;
@@ -1436,14 +1532,14 @@ namespace
                        indices.Get(), 0, &indexSource);
                }
                const D3D11_BOX argsSource{outputArgsByteOffset,0,0,
-                    outputArgsByteOffset+10u*sizeof(std::uint32_t),1,1};
+                    outputArgsByteOffset+11u*sizeof(std::uint32_t),1,1};
                context->CopySubresourceRegion(argsStaging.Get(),0,0,0,0,args.Get(),0,&argsSource);
                std::uint32_t connectorBase=0;
-                if(farField&&farField->connectorVertexCount)
+                if(hasConnectorOutput)
                {
                    connectorBase=pooledChunk?outputBaseVertex/farMaxVerticesPerTile*FarPanelVertexCapacity:0;
-                    const UINT vertexBytes=farField->connectorVertexCount*sizeof(GpuVertex);
-                    const UINT indexBytes=farField->connectorVertexCount*sizeof(std::uint32_t);
+                    const UINT vertexBytes=connectorPanelCapacity*sizeof(GpuVertex);
+                    const UINT indexBytes=connectorPanelCapacity*sizeof(std::uint32_t);
                    const D3D11_BOX connectorVertexSource{connectorBase*sizeof(GpuVertex),0,0,
                        connectorBase*sizeof(GpuVertex)+vertexBytes,1,1};
                    const D3D11_BOX connectorIndexSource{connectorBase*sizeof(std::uint32_t),0,0,
@@ -1490,7 +1586,8 @@ namespace
                    if(profileGpuKernels&&!disjoint.Disjoint&&disjoint.Frequency)
                        for(std::size_t stage=0;stage<result.kernelMs.size();++stage)
                        {
-                            if(stage==5&&farField->connectorVertexCount==0)continue;
+                            if((stage==5&&farField->connectorVertexCount==0)
+                                ||(stage>=6&&farField->visibleSkirtFaceMask==0))continue;
                            const UINT64 stageStart=ReadQuery<UINT64>(kernelStartQueries[stage].Get(),
                                "GetData(kernel stage start)");
                            const UINT64 stageEnd=ReadQuery<UINT64>(kernelEndQueries[stage].Get(),
@@ -1505,18 +1602,23 @@ namespace
                    static_cast<std::uint32_t*>(mapped.pData) + cellCount);
                context->Unmap(caseStaging.Get(), 0);
                Check(context->Map(argsStaging.Get(), 0, D3D11_MAP_READ, 0, &mapped), "Map(args)");
-                std::array<std::uint32_t,10> allArgs{};
-                std::copy_n(static_cast<const std::uint32_t*>(mapped.pData),10,allArgs.begin());
+                std::array<std::uint32_t,11> allArgs{};
+                std::copy_n(static_cast<const std::uint32_t*>(mapped.pData),11,allArgs.begin());
                std::copy_n(allArgs.begin(),5,result.drawArgs.begin());
+                std::copy_n(allArgs.begin()+5,6,result.connectorDrawArgs.begin());
                context->Unmap(argsStaging.Get(), 0);
                result.indexedIndirectArgsValid = (useIndirectOutputCount
                    ?result.drawArgs[0]<=maxVertices
                    :result.drawArgs[0] == tile.expectedVertexCount)
                    && result.drawArgs[1] == 1 && result.drawArgs[2] == outputBaseVertex
                    && result.drawArgs[3] == 0 && result.drawArgs[4] == 0;
-                result.indexedIndirectArgsValid &= allArgs[5]==(farField?farField->connectorVertexCount:0)
-                    &&allArgs[6]==(farField&&farField->connectorVertexCount?1u:0u)
-                    &&allArgs[7]==connectorBase&&allArgs[8]==connectorBase&&allArgs[9]==0;
+                const std::uint32_t connectorOutputCount=std::min<std::uint32_t>(allArgs[5],
+                    connectorPanelCapacity);
+                result.indexedIndirectArgsValid &= allArgs[5]>=
+                    (farField?farField->connectorVertexCount:0)
+                    &&allArgs[6]==(connectorOutputCount>0?1u:0u)
+                    &&allArgs[7]==connectorBase&&allArgs[8]==0&&allArgs[9]==0
+                    &&allArgs[10]==0;
                const std::uint32_t outputCount = std::min<std::uint32_t>(result.drawArgs[0],
                    static_cast<std::uint32_t>(maxVertices));
                // The staging buffer contains the GPU output. Only retain the prefix named by
@@ -1531,18 +1633,22 @@ namespace
                for (std::uint32_t i = 0; i < outputCount; ++i)
                    result.sequentialIndices &= indexData[i] == outputBaseVertex+i;
                context->Unmap(indexStaging.Get(), 0);
-                if(farField&&farField->connectorVertexCount)
+                if(hasConnectorOutput&&connectorOutputCount>0)
                {
                    Check(context->Map(connectorVertexStaging.Get(),0,D3D11_MAP_READ,0,&mapped),
-                        "Map(far connector panel vertices)");
+                        "Map(far connector and skirt vertices)");
                    const auto* connectorData=static_cast<const GpuVertex*>(mapped.pData);
+                    const std::uint32_t panelOutputCount=std::min<std::uint32_t>(
+                        farField->connectorVertexCount,connectorOutputCount);
                    result.vertices.insert(result.vertices.end(),connectorData,
-                        connectorData+farField->connectorVertexCount);
+                        connectorData+panelOutputCount);
+                    result.skirtVertices.assign(connectorData+panelOutputCount,
+                        connectorData+connectorOutputCount);
                    context->Unmap(connectorVertexStaging.Get(),0);
                    Check(context->Map(connectorIndexStaging.Get(),0,D3D11_MAP_READ,0,&mapped),
-                        "Map(far connector panel indices)");
+                        "Map(far connector and skirt indices)");
                    const auto* connectorIndexData=static_cast<const std::uint32_t*>(mapped.pData);
-                    for(std::uint32_t i=0;i<farField->connectorVertexCount;++i)
+                    for(std::uint32_t i=0;i<connectorOutputCount;++i)
                        result.sequentialIndices &= connectorIndexData[i]==connectorBase+i;
                    context->Unmap(connectorIndexStaging.Get(),0);
                }
@@ -1618,6 +1724,7 @@ namespace
        TileKey key;
        FarFieldTileInput input;
        double exportPreparationMs=0.0;
+        std::uint8_t skirtFaceMask=0;
    };

    struct SpatialKey
@@ -1892,6 +1999,38 @@ namespace
        return result;
    }

+    std::uint8_t ComputeVisibleFarSkirtFaceMask(const TileKey& tile,
+        const std::unordered_set<TileKey,TileKeyHash>& desired)
+    {
+        const std::int64_t extent=static_cast<std::int64_t>(32)<<tile.lod;
+        const std::int64_t origin[3]={tile.coordinate.x*extent,
+            tile.coordinate.y*extent,tile.coordinate.z*extent};
+        const auto floorDiv=[](std::int64_t value,std::int64_t divisor)
+        {
+            std::int64_t quotient=value/divisor;
+            if(value%divisor<0)--quotient;
+            return static_cast<std::int32_t>(quotient);
+        };
+        std::uint8_t mask=0;
+        for(int axis=0;axis<3;++axis)
+        for(int side=0;side<2;++side)
+        {
+            std::int64_t probe[3]={origin[0]+extent/2,origin[1]+extent/2,origin[2]+extent/2};
+            probe[axis]=side?origin[axis]+extent:origin[axis]-1;
+            int neighborLod=-1;
+            for(int lod=0;lod<=8;++lod)
+            {
+                const std::int64_t neighborExtent=static_cast<std::int64_t>(32)<<lod;
+                const TileKey candidate{{floorDiv(probe[0],neighborExtent),
+                    floorDiv(probe[1],neighborExtent),floorDiv(probe[2],neighborExtent)},lod};
+                if(desired.contains(candidate)){neighborLod=lod;break;}
+            }
+            if(neighborLod<0||neighborLod!=tile.lod)
+                mask|=static_cast<std::uint8_t>(1u<<(axis*2+side));
+        }
+        return mask;
+    }
+
    void PrintProcessMemory(const char* stage)
    {
        PROCESS_MEMORY_COUNTERS_EX memory{};
@@ -1906,6 +2045,185 @@ namespace
        }
    }

+    struct SkirtAuditResult
+    {
+        std::uint64_t cpuTriangles=0,gpuTriangles=0,topologyDifferenceTriangles=0;
+        std::uint64_t quantizedKeyDifferenceTriangles=0;
+        double meanVertexDistanceCm=0.0;
+        float maxVertexDistanceCm=0.0f;
+        float vertexBoundCm=0.15f;
+        bool pass=false;
+    };
+
+    SkirtAuditResult AuditGpuSkirts(World& world,const TileMesh& cpuMesh,
+        const std::vector<GpuVertex>& gpuSkirtVertices,std::uint8_t visibleFaceMask)
+    {
+        using FQuantizedVertex=std::array<std::int64_t,3>;
+        using FTriangleKey=std::array<std::int64_t,9>;
+        std::map<FTriangleKey,std::uint32_t> cpuTriangles,gpuTriangles;
+        std::vector<GpuVertex> cpuVertices,gpuVertices;
+        std::vector<std::array<GpuVertex,3>> cpuGeometry,gpuGeometry;
+        const auto makeVertex=[](double x,double y,double z,double nx,double ny,double nz)
+        {
+            GpuVertex vertex{};
+            vertex.position[0]=static_cast<float>(x);
+            vertex.position[1]=static_cast<float>(y);
+            vertex.position[2]=static_cast<float>(z);
+            vertex.normal[0]=static_cast<float>(nx);
+            vertex.normal[1]=static_cast<float>(ny);
+            vertex.normal[2]=static_cast<float>(nz);
+            return vertex;
+        };
+        const auto triangleKey=[](const GpuVertex* vertices)
+        {
+            std::array<FQuantizedVertex,3> points{};
+            for(std::size_t i=0;i<3;++i)
+                points[i]={static_cast<std::int64_t>(std::llround(vertices[i].position[0])),
+                    static_cast<std::int64_t>(std::llround(vertices[i].position[1])),
+                    static_cast<std::int64_t>(std::llround(vertices[i].position[2]))};
+            std::sort(points.begin(),points.end());
+            FTriangleKey key{};std::size_t out=0;
+            for(const auto& point:points)for(const auto coordinate:point)key[out++]=coordinate;
+            return key;
+        };
+        const std::size_t cpuTriangleCount=cpuMesh.indices.size()/3;
+        for(std::size_t triangle=0;triangle<cpuTriangleCount;++triangle)
+        {
+            const std::uint8_t faceMask=triangle<cpuMesh.skirtFaceMasks.size()
+                ?cpuMesh.skirtFaceMasks[triangle]:0;
+            if((faceMask&visibleFaceMask)==0)continue;
+            GpuVertex points[3]{};
+            bool validTriangle=true;
+            for(std::size_t corner=0;corner<3;++corner)
+            {
+                const std::uint32_t index=cpuMesh.indices[triangle*3+corner];
+                if(index>=cpuMesh.vertices.size()){validTriangle=false;break;}
+                const MeshVertex& source=cpuMesh.vertices[index];
+                points[corner]=makeVertex(source.position.x,source.position.y,source.position.z,
+                    source.normal.x,source.normal.y,source.normal.z);
+                cpuVertices.push_back(points[corner]);
+            }
+            if(!validTriangle)continue;
+            cpuGeometry.push_back({points[0],points[1],points[2]});
+            ++cpuTriangles[triangleKey(points)];
+        }
+        for(std::size_t triangle=0;triangle+2<gpuSkirtVertices.size();triangle+=3)
+        {
+            const GpuVertex points[3]={gpuSkirtVertices[triangle],gpuSkirtVertices[triangle+1],
+                gpuSkirtVertices[triangle+2]};
+            gpuGeometry.push_back({points[0],points[1],points[2]});
+            ++gpuTriangles[triangleKey(points)];
+            gpuVertices.insert(gpuVertices.end(),std::begin(points),std::end(points));
+        }
+        SkirtAuditResult result;
+        result.cpuTriangles=cpuVertices.size()/3;
+        result.gpuTriangles=gpuVertices.size()/3;
+        auto cpuIt=cpuTriangles.begin(),gpuIt=gpuTriangles.begin();
+        while(cpuIt!=cpuTriangles.end()||gpuIt!=gpuTriangles.end())
+        {
+            if(gpuIt==gpuTriangles.end()||(cpuIt!=cpuTriangles.end()&&cpuIt->first<gpuIt->first))
+            {result.quantizedKeyDifferenceTriangles+=cpuIt->second;++cpuIt;}
+            else if(cpuIt==cpuTriangles.end()||gpuIt->first<cpuIt->first)
+            {result.quantizedKeyDifferenceTriangles+=gpuIt->second;++gpuIt;}
+            else
+            {
+                result.quantizedKeyDifferenceTriangles+=cpuIt->second>gpuIt->second
+                    ?cpuIt->second-gpuIt->second:gpuIt->second-cpuIt->second;
+                ++cpuIt;++gpuIt;
+            }
+        }
+        const auto vertexDistance=[](const GpuVertex& a,const GpuVertex& b)
+        {
+            const double dx=static_cast<double>(a.position[0])-b.position[0];
+            const double dy=static_cast<double>(a.position[1])-b.position[1];
+            const double dz=static_cast<double>(a.position[2])-b.position[2];
+            return std::sqrt(dx*dx+dy*dy+dz*dz);
+        };
+        constexpr std::uint8_t permutations[6][3]={
+            {0,1,2},{0,2,1},{1,0,2},{1,2,0},{2,0,1},{2,1,0}};
+        std::vector<bool> matched(gpuGeometry.size(),false);
+        for(const auto& cpuTriangle:cpuGeometry)
+        {
+            std::size_t best=gpuGeometry.size();double bestError=std::numeric_limits<double>::infinity();
+            for(std::size_t candidate=0;candidate<gpuGeometry.size();++candidate)
+            {
+                if(matched[candidate])continue;
+                for(const auto& permutation:permutations)
+                {
+                    double maximum=0.0,sum=0.0;
+                    for(std::size_t corner=0;corner<3;++corner)
+                    {
+                        const double distance=vertexDistance(cpuTriangle[corner],
+                            gpuGeometry[candidate][permutation[corner]]);
+                        maximum=std::max(maximum,distance);sum+=distance;
+                    }
+                    if(maximum<=0.15&&sum<bestError)
+                    {best=candidate;bestError=sum;}
+                }
+            }
+            if(best==gpuGeometry.size())++result.topologyDifferenceTriangles;
+            else matched[best]=true;
+        }
+        result.topologyDifferenceTriangles+=static_cast<std::uint64_t>(
+            std::count(matched.begin(),matched.end(),false));
+        const DistanceStats cpuToGpu=NearestDistance(cpuVertices,gpuVertices);
+        const DistanceStats gpuToCpu=NearestDistance(gpuVertices,cpuVertices);
+        const std::uint64_t comparisonCount=cpuToGpu.count+gpuToCpu.count;
+        result.meanVertexDistanceCm=comparisonCount
+            ?(cpuToGpu.meanCm*cpuToGpu.count+gpuToCpu.meanCm*gpuToCpu.count)/comparisonCount:0.0;
+        result.maxVertexDistanceCm=std::max(cpuToGpu.maxCm,gpuToCpu.maxCm);
+        if(result.maxVertexDistanceCm>1.0f)
+        {
+            const auto normalAngleDegrees=[](const GpuVertex& a,const GpuVertex& b)
+            {
+                const double al=std::sqrt(double(a.normal[0])*a.normal[0]
+                    +double(a.normal[1])*a.normal[1]+double(a.normal[2])*a.normal[2]);
+                const double bl=std::sqrt(double(b.normal[0])*b.normal[0]
+                    +double(b.normal[1])*b.normal[1]+double(b.normal[2])*b.normal[2]);
+                if(al<=1.0e-9||bl<=1.0e-9)return 180.0;
+                const double dot=(double(a.normal[0])*b.normal[0]
+                    +double(a.normal[1])*b.normal[1]+double(a.normal[2])*b.normal[2])/(al*bl);
+                return std::acos(std::clamp(dot,-1.0,1.0))*57.29577951308232;
+            };
+            const auto fieldDensity=[&](const GpuVertex& vertex)
+            {
+                return world.densityAt({static_cast<double>(vertex.position[0])/25.0,
+                    static_cast<double>(vertex.position[1])/25.0,
+                    static_cast<double>(vertex.position[2])/25.0});
+            };
+            double maxCpuToGpu=0.0,maxGpuToCpu=0.0;
+            const GpuVertex* cpuWorst=nullptr;const GpuVertex* gpuNearest=nullptr;
+            const GpuVertex* gpuWorst=nullptr;const GpuVertex* cpuNearest=nullptr;
+            for(const GpuVertex& cpu:cpuVertices)
+            {
+                double best=std::numeric_limits<double>::infinity();const GpuVertex* nearest=nullptr;
+                for(const GpuVertex& gpu:gpuVertices)
+                {const double distance=vertexDistance(cpu,gpu);if(distance<best){best=distance;nearest=&gpu;}}
+                if(best>maxCpuToGpu){maxCpuToGpu=best;cpuWorst=&cpu;gpuNearest=nearest;}
+            }
+            for(const GpuVertex& gpu:gpuVertices)
+            {
+                double best=std::numeric_limits<double>::infinity();const GpuVertex* nearest=nullptr;
+                for(const GpuVertex& cpu:cpuVertices)
+                {const double distance=vertexDistance(gpu,cpu);if(distance<best){best=distance;nearest=&cpu;}}
+                if(best>maxGpuToCpu){maxGpuToCpu=best;gpuWorst=&gpu;cpuNearest=nearest;}
+            }
+            if(cpuWorst&&gpuNearest&&gpuWorst&&cpuNearest)
+                std::printf("[GpuSkirtAuditWorst] cpu_to_gpu_cm=%.6f cpu=(%.5f,%.5f,%.5f) cpu_field=%.9g gpu_near=(%.5f,%.5f,%.5f) gpu_near_cpu_field=%.9g normal_angle_deg=%.6f gpu_to_cpu_cm=%.6f gpu=(%.5f,%.5f,%.5f) gpu_cpu_field=%.9g cpu_near=(%.5f,%.5f,%.5f) cpu_near_field=%.9g normal_angle_deg=%.6f\n",
+                    maxCpuToGpu,cpuWorst->position[0],cpuWorst->position[1],cpuWorst->position[2],
+                    fieldDensity(*cpuWorst),gpuNearest->position[0],gpuNearest->position[1],
+                    gpuNearest->position[2],fieldDensity(*gpuNearest),
+                    normalAngleDegrees(*cpuWorst,*gpuNearest),maxGpuToCpu,gpuWorst->position[0],
+                    gpuWorst->position[1],gpuWorst->position[2],fieldDensity(*gpuWorst),
+                    cpuNearest->position[0],cpuNearest->position[1],cpuNearest->position[2],
+                    fieldDensity(*cpuNearest),
+                    normalAngleDegrees(*gpuWorst,*cpuNearest));
+        }
+        result.pass=result.cpuTriangles==result.gpuTriangles
+            &&result.topologyDifferenceTriangles==0&&result.maxVertexDistanceCm<=result.vertexBoundCm;
+        return result;
+    }
+
    bool ReportTile(World& world,const TileData& tile, const TileMesh& cpuMesh,
                    const GpuResult& hwFirst, const GpuResult& hwSecond,
                    const GpuResult& warp,const FarFieldTileInput& farInput)
@@ -1928,6 +2246,8 @@ namespace
            repeatCaseDiff += hwFirst.cases[i] != hwSecond.cases[i];
        for (std::size_t i = 0; i < hwFirst.drawArgs.size(); ++i)
            repeatArgsDiff += hwFirst.drawArgs[i] != hwSecond.drawArgs[i];
+        for(std::size_t i=0;i<hwFirst.connectorDrawArgs.size();++i)
+            repeatArgsDiff+=hwFirst.connectorDrawArgs[i]!=hwSecond.connectorDrawArgs[i];
        const std::size_t repeatCount = std::min(hwFirst.vertices.size(), hwSecond.vertices.size());
        for (std::size_t i = 0; i < repeatCount; ++i)
        {
@@ -1935,6 +2255,20 @@ namespace
            const std::uint32_t* b = reinterpret_cast<const std::uint32_t*>(&hwSecond.vertices[i]);
            for (int c = 0; c < 8; ++c) repeatBitDiff += a[c] != b[c];
        }
+        repeatBitDiff+=hwFirst.skirtVertices.size()!=hwSecond.skirtVertices.size();
+        for(std::size_t i=0;i<std::min(hwFirst.skirtVertices.size(),hwSecond.skirtVertices.size());++i)
+        {
+            const std::uint32_t* a=reinterpret_cast<const std::uint32_t*>(&hwFirst.skirtVertices[i]);
+            const std::uint32_t* b=reinterpret_cast<const std::uint32_t*>(&hwSecond.skirtVertices[i]);
+            for(int c=0;c<8;++c)repeatBitDiff+=a[c]!=b[c];
+        }
+        warpBitDiff+=hwFirst.skirtVertices.size()!=warp.skirtVertices.size();
+        for(std::size_t i=0;i<std::min(hwFirst.skirtVertices.size(),warp.skirtVertices.size());++i)
+        {
+            const std::uint32_t* a=reinterpret_cast<const std::uint32_t*>(&hwFirst.skirtVertices[i]);
+            const std::uint32_t* b=reinterpret_cast<const std::uint32_t*>(&warp.skirtVertices[i]);
+            for(int c=0;c<8;++c)warpBitDiff+=a[c]!=b[c];
+        }
        const std::size_t paired = std::min(hwFirst.vertices.size(), warp.vertices.size());
        for (std::size_t i = 0; i < paired; ++i)
        {
@@ -2132,6 +2466,34 @@ namespace
            static_cast<unsigned long long>(fieldSignDiff),
            static_cast<unsigned long long>(fieldCaseLatticeSignDiff),
            static_cast<unsigned long long>(fieldHaloSignDiff));
+        SkirtAuditResult skirtAudit{};
+        SkirtAuditResult warpSkirtAudit{};
+        bool skirtAuditPass=hwFirst.skirtVertices.empty()&&warp.skirtVertices.empty();
+        if(farInput.visibleSkirtFaceMask!=0)
+        {
+            skirtAudit=AuditGpuSkirts(world,cpuMesh,hwFirst.skirtVertices,farInput.visibleSkirtFaceMask);
+            warpSkirtAudit=AuditGpuSkirts(world,cpuMesh,warp.skirtVertices,farInput.visibleSkirtFaceMask);
+            skirtAuditPass=skirtAudit.pass&&warpSkirtAudit.pass;
+            std::printf("[GpuSkirtAudit] tile=(%d,%d,%d,lod%d) face_mask=0x%02x cpu_triangles=%llu gpu_triangles=%llu topology_diff_triangles_at_0.15cm=%llu quantized_key_diff_triangles_1cm=%llu vertex_mean_cm=%.8g vertex_max_cm=%.8g vertex_bound_cm=%.2f connector_output_vertices=%zu skirt_vertices=%zu overflow=%u status=%s\n",
+                tile.key.coordinate.x,tile.key.coordinate.y,tile.key.coordinate.z,tile.key.lod,
+                static_cast<unsigned>(farInput.visibleSkirtFaceMask),
+                static_cast<unsigned long long>(skirtAudit.cpuTriangles),
+                static_cast<unsigned long long>(skirtAudit.gpuTriangles),
+                static_cast<unsigned long long>(skirtAudit.topologyDifferenceTriangles),
+                static_cast<unsigned long long>(skirtAudit.quantizedKeyDifferenceTriangles),
+                skirtAudit.meanVertexDistanceCm,skirtAudit.maxVertexDistanceCm,
+                skirtAudit.vertexBoundCm,hwFirst.vertices.size()-tile.expectedVertexCount,
+                hwFirst.skirtVertices.size(),hwFirst.connectorDrawArgs[5],skirtAudit.pass?"PASS":"FAIL");
+            std::printf("[WarpSkirtAudit] tile=(%d,%d,%d,lod%d) face_mask=0x%02x cpu_triangles=%llu warp_triangles=%llu topology_diff_triangles_at_0.15cm=%llu quantized_key_diff_triangles_1cm=%llu vertex_mean_cm=%.8g vertex_max_cm=%.8g vertex_bound_cm=%.2f overflow=%u status=%s\n",
+                tile.key.coordinate.x,tile.key.coordinate.y,tile.key.coordinate.z,tile.key.lod,
+                static_cast<unsigned>(farInput.visibleSkirtFaceMask),
+                static_cast<unsigned long long>(warpSkirtAudit.cpuTriangles),
+                static_cast<unsigned long long>(warpSkirtAudit.gpuTriangles),
+                static_cast<unsigned long long>(warpSkirtAudit.topologyDifferenceTriangles),
+                static_cast<unsigned long long>(warpSkirtAudit.quantizedKeyDifferenceTriangles),
+                warpSkirtAudit.meanVertexDistanceCm,warpSkirtAudit.maxVertexDistanceCm,
+                warpSkirtAudit.vertexBoundCm,warp.connectorDrawArgs[5],warpSkirtAudit.pass?"PASS":"FAIL");
+        }
        const std::size_t expectedGpuOutputCount=static_cast<std::size_t>(tile.expectedVertexCount)
            +farInput.connectorVertexCount;
        const bool outputCountMatch = hwFirst.vertices.size() == expectedGpuOutputCount
@@ -2140,12 +2502,14 @@ namespace
        const bool vertexAgreement=maximum<=vertexBoundCm&&cpuWarpMaximum<=vertexBoundCm&&warpMaximum<=vertexBoundCm;
        const bool repeatedOutputsIdentical = repeatBitDiff == 0 && repeatCaseDiff == 0
            && repeatArgsDiff == 0 && hwFirst.vertices.size() == hwSecond.vertices.size()
+            &&hwFirst.skirtVertices.size()==hwSecond.skirtVertices.size()
            && hwFirst.sequentialIndices && hwSecond.sequentialIndices
            && hwFirst.indexedIndirectArgsValid && hwSecond.indexedIndirectArgsValid;
        const bool status = hwTopologyDiff == 0 && warpTopologyDiff == 0
            &&fieldCount==tile.density.size()
            &&warpFieldCount==tile.density.size()
            && repeatedOutputsIdentical && outputCountMatch
+            &&skirtAuditPass
            &&protectedAirSignDiffHw==0&&protectedAirSignDiffWarp==0
            &&protectedSolidSignDiffHw==0&&protectedSolidSignDiffWarp==0
            &&thinFeatureDiffGpu==0&&thinFeatureDiffWarp==0
@@ -2190,6 +2554,7 @@ namespace
        bool wholeFar=false;
        bool wholeNear=false;
        bool profileGpuKernels=false;
+        bool skirtAuditOnly=false;
        int clipRadiusOverride=-1;
        int tileLimit=0;
        bool hasCameraOverride=false;
@@ -2199,6 +2564,7 @@ namespace
            if(std::string_view(argv[i])=="--whole-far")wholeFar=true;
            if(std::string_view(argv[i])=="--whole-near")wholeNear=true;
            if(std::string_view(argv[i])=="--profile-kernels")profileGpuKernels=true;
+            if(std::string_view(argv[i])=="--skirt-audit")skirtAuditOnly=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()));
@@ -2276,13 +2642,32 @@ 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,wholeNear);
+        std::vector<FarSelection> selected = skirtAuditOnly
+            ?std::vector<FarSelection>{}:SelectFarTiles(world, config, camera,wholeFar,wholeNear);
+        if(skirtAuditOnly)
+        {
+            if(!config.generateSkirts)
+                throw std::runtime_error("--skirt-audit requires the owner's skirt option to be enabled");
+            const std::unordered_set<TileKey,TileKeyHash> desiredSet(desired.begin(),desired.end());
+            selected.clear();
+            for(const TileKey& key:desired)
+            {
+                if(key.lod<std::max(config.fullResolutionLodCount,2)||key.lod>config.maxLod
+                    ||!world.isFarFieldTileEligible(key))continue;
+                const std::uint8_t mask=ComputeVisibleFarSkirtFaceMask(key,desiredSet);
+                if(mask==0)continue;
+                selected.push_back({key,{},0.0,mask});
+            }
+            if(selected.empty())throw std::runtime_error("no desired far tile had exposed CPU skirt triangles");
+            std::printf("[GpuSkirtAuditSelection] candidate_tiles=%zu mode=all-visible-far-faces\n",
+                selected.size());
+        }
        if (selected.empty()) throw std::runtime_error("no far-field tiles were selected");
        bool allPass = noiseCpuGpu.differingBits == 0 && noiseRepeat.differingBits == 0
            && noiseWarp.differingBits == 0;
        std::uint64_t testedTiles=0;double cpuTileWallSum=0.0,cpuFieldPrepSum=0.0,cpuExportPrepSum=0.0;
        double gpuFieldSum=0.0,gpuTotalSum=0.0,warpFieldSum=0.0;
-        std::array<double,6> gpuKernelStageSum{};
+        std::array<double,9> gpuKernelStageSum{};
        std::uint64_t gpuFieldDistanceEvaluations=0,gpuFieldSamples=0;
        std::uint32_t gpuFieldMaxDistanceEvaluationsPerSample=0;
        std::uint64_t nearTransitionSkipped=0,topologyDiffGpuTotal=0,topologyDiffWarpTotal=0;
@@ -2296,7 +2681,7 @@ namespace
            FarFieldTileInput farInputStorage;
            const FarFieldTileInput* farInput=&selection.input;
            double exportMs=selection.exportPreparationMs;
-            if(wholeFar||wholeNear)
+            if((wholeFar||wholeNear)&&!skirtAuditOnly)
            {
                const auto exportStart=Clock::now();
                auto exported=wholeNear?world.exportGpuFieldTile(key,error)
@@ -2322,6 +2707,25 @@ namespace
            auto cpuMesh = world.generateTile(key, options, error);
            const double cpuWallMs = std::chrono::duration<double, std::milli>(Clock::now() - cpuStart).count();
            if (!cpuMesh) throw std::runtime_error("CPU reference tile failed: " + error);
+            if(skirtAuditOnly)
+            {
+                std::uint64_t exposedTriangles=0;
+                for(const std::uint8_t faces:cpuMesh->skirtFaceMasks)
+                    exposedTriangles+=(faces&selection.skirtFaceMask)!=0;
+                const auto exportStart=Clock::now();
+                auto exported=world.exportFarFieldTile(key,error);
+                exportMs=std::chrono::duration<double,std::milli>(Clock::now()-exportStart).count();
+                if(!exported)throw std::runtime_error("skirt audit far field export failed: "+error);
+                farInputStorage=std::move(*exported);
+                farInputStorage.visibleSkirtFaceMask=selection.skirtFaceMask;
+                farInputStorage.skirtDepthVoxels=static_cast<float>(std::max(1,config.skirtCells)
+                    *farInputStorage.step);
+                farInput=&farInputStorage;
+                std::printf("[GpuSkirtAuditSelection] key=(%d,%d,%d,lod%d) face_mask=0x%02x cpu_exposed_triangles=%llu\n",
+                    key.coordinate.x,key.coordinate.y,key.coordinate.z,key.lod,
+                    static_cast<unsigned>(selection.skirtFaceMask),
+                    static_cast<unsigned long long>(exposedTriangles));
+            }
            const auto inputStart = Clock::now();
            TileData tile = MakeTileData(world, key, config, *cpuMesh, true,farInput);
            const double inputWallMs = std::chrono::duration<double, std::milli>(Clock::now() - inputStart).count();
@@ -2359,11 +2763,14 @@ namespace
                &&gpu1.vertices.size() == static_cast<std::size_t>(tile.expectedVertexCount)+farInput->connectorVertexCount
                && gpu2.vertices.size() == gpu1.vertices.size()
                && warpResult.vertices.size() == static_cast<std::size_t>(tile.expectedVertexCount)+farInput->connectorVertexCount
+                &&gpu1.skirtVertices==gpu2.skirtVertices
+                &&warpResult.skirtVertices.size()==gpu1.skirtVertices.size()
                && gpu1.sequentialIndices && gpu2.sequentialIndices && warpResult.sequentialIndices
                && gpu1.indexedIndirectArgsValid && gpu2.indexedIndirectArgsValid
                && warpResult.indexedIndirectArgsValid
                && gpu1.cases == tile.expectedCases && gpu2.cases == tile.expectedCases
                && warpResult.cases == tile.expectedCases && gpu1.drawArgs == gpu2.drawArgs
+                &&gpu1.connectorDrawArgs==gpu2.connectorDrawArgs
                && agreementPass;
            allPass &= tilePass;
            ++testedTiles;cpuTileWallSum+=cpuWallMs;cpuFieldPrepSum+=inputWallMs;
@@ -2403,9 +2810,10 @@ namespace
            static_cast<unsigned long long>(protectedSolidOverlaySamples));
        if(profileGpuKernels)
        {
-            std::printf("[GpuKernelProfile] tiles=%llu field_ms=%.6f classify_ms=%.6f scan_ms=%.6f emit_ms=%.6f connector_rings_ms=%.6f connector_panels_ms=%.6f field_shape_distance_evaluations=%llu field_lattice_samples=%llu average_shape_distance_evaluations_per_lattice_sample=%.4f max_shape_distance_evaluations_per_lattice_sample=%u\n",
+            std::printf("[GpuKernelProfile] tiles=%llu field_ms=%.6f classify_ms=%.6f scan_ms=%.6f emit_ms=%.6f connector_rings_ms=%.6f connector_panels_ms=%.6f skirt_classify_ms=%.6f skirt_scan_ms=%.6f skirt_emit_ms=%.6f field_shape_distance_evaluations=%llu field_lattice_samples=%llu average_shape_distance_evaluations_per_lattice_sample=%.4f max_shape_distance_evaluations_per_lattice_sample=%u\n",
                static_cast<unsigned long long>(testedTiles),gpuKernelStageSum[0],gpuKernelStageSum[1],
                gpuKernelStageSum[2],gpuKernelStageSum[3],gpuKernelStageSum[4],gpuKernelStageSum[5],
+                gpuKernelStageSum[6],gpuKernelStageSum[7],gpuKernelStageSum[8],
                static_cast<unsigned long long>(gpuFieldDistanceEvaluations),
                static_cast<unsigned long long>(gpuFieldSamples),
                gpuFieldSamples?static_cast<double>(gpuFieldDistanceEvaluations)/gpuFieldSamples:0.0,
@@ -2544,7 +2952,8 @@ namespace VoxelForgeGpuFarViewer
                auditReadback!=nullptr,auditReadback!=nullptr,outputSlot,true,captureGpuTiming,batchTimingMode);
            if(auditReadback)
            {
-                auditReadback->vertices.resize(result.vertices.size());
+                const std::size_t completeGpuVertexCount=result.vertices.size()+result.skirtVertices.size();
+                auditReadback->vertices.resize(completeGpuVertexCount);
                for(std::size_t i=0;i<result.vertices.size();++i)
                {
                    for(int c=0;c<3;++c)
@@ -2553,7 +2962,16 @@ namespace VoxelForgeGpuFarViewer
                        auditReadback->vertices[i].normal[c]=result.vertices[i].normal[c];
                    }
                }
-                auditReadback->indices.resize(result.vertices.size());
+                for(std::size_t i=0;i<result.skirtVertices.size();++i)
+                {
+                    const GpuVertex& source=result.skirtVertices[i];
+                    for(int c=0;c<3;++c)
+                    {
+                        auditReadback->vertices[result.vertices.size()+i].position[c]=source.position[c];
+                        auditReadback->vertices[result.vertices.size()+i].normal[c]=source.normal[c];
+                    }
+                }
+                auditReadback->indices.resize(completeGpuVertexCount);
                std::iota(auditReadback->indices.begin(),auditReadback->indices.end(),0u);
            }
            if(completionEvent)
@@ -0,0 +1,73 @@
+[CmdletBinding()]
+param(
+    [string]$ExecutableName='VoxelForgeGpuFieldStep8SkirtViewerC.exe',
+    [string]$CapturePrefix='step8'
+)
+
+$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($CapturePrefix -notmatch '^[A-Za-z0-9_-]+$')
+{throw 'CapturePrefix must contain only letters, digits, underscores, or hyphens.'}
+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="${CapturePrefix}_${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 "Full-residency capture or 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
+        }
+    }
+}
@@ -0,0 +1,56 @@
+[CmdletBinding()]
+param([string]$ExecutableName='VoxelForgeGpuFieldStep8SkirtViewerC.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='step8_interleaved_cpu_A'},
+    @{Mode='gpu';Name='step8_interleaved_gpu_A'},
+    @{Mode='gpu';Name='step8_interleaved_gpu_B'},
+    @{Mode='cpu';Name='step8_interleaved_cpu_B'},
+    @{Mode='cpu';Name='step8_interleaved_cpu_C'},
+    @{Mode='gpu';Name='step8_interleaved_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 "Full 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
+    }
+}
@@ -0,0 +1,70 @@
+[CmdletBinding()]
+param(
+    [string]$ExecutableName='VoxelForgeGpuFieldStep7SkirtSeamVerify.exe',
+    [int]$IdleSeconds=60
+)
+
+$ErrorActionPreference='Stop'
+$HostProcess=Get-Process -Id $PID
+$HostProcess.PriorityClass='BelowNormal'
+$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$BinRoot=Join-Path $StandaloneRoot 'bin'
+$ExecutablePath=Join-Path $BinRoot $ExecutableName
+$RenamedPath=Join-Path $BinRoot 'VoxelForgeGpuFieldStep8PresentRenameControl.exe'
+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"}
+Copy-Item -LiteralPath $ExecutablePath -Destination $RenamedPath -Force
+
+$Sequence=@(
+    @{Mode='gpu';Name='step8_present_same_A';Path=$ExecutablePath;Delay=0},
+    @{Mode='gpu';Name='step8_present_same_B';Path=$ExecutablePath;Delay=0},
+    @{Mode='cpu';Name='step8_present_cpu_control';Path=$ExecutablePath;Delay=0},
+    @{Mode='gpu';Name='step8_present_after_cpu';Path=$ExecutablePath;Delay=0},
+    @{Mode='gpu';Name='step8_present_after_idle';Path=$ExecutablePath;Delay=$IdleSeconds},
+    @{Mode='gpu';Name='step8_present_renamed_exe';Path=$RenamedPath;Delay=0}
+)
+foreach($Run in $Sequence)
+{
+    if($Run.Delay-gt0)
+    {
+        Write-Output "CONTROL_DELAY seconds=$($Run.Delay) before=$($Run.Name)"
+        Start-Sleep -Seconds $Run.Delay
+    }
+    $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 "PRESENT_CONTROL command=`"$($Run.Path)`" $($Arguments -join ' ')"
+    for($Attempt=0;$Attempt-lt2;$Attempt++)
+    {
+        if($Attempt-gt0)
+        {Write-Output '0xC0000142 encountered; waiting 60 seconds before the single permitted retry.';Start-Sleep -Seconds 60}
+        Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
+        $Child=Start-Process -FilePath $Run.Path -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
+            -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
+        $PeakPrivateMemory=0L;$OverMemoryLimit=$false
+        while(-not$Child.HasExited)
+        {
+            $Child.Refresh();if($Child.HasExited){break}
+            $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+            $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()
+        $ExitCode=0
+        if($null-ne$Child.ExitCode){$ExitCode=[int]$Child.ExitCode}
+        if($ExitCode-eq-1073741502-and$Attempt-eq0){continue}
+        if($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=$ExitCode"}
+        $PresentLine=Select-String -Path $StdoutPath -Pattern '^\[PresentProfile\] ordinal=1 '
+        Write-Output "CONTROL_EXIT mode=$($Run.Mode) name=$($Run.Name) pid=$($Child.Id) priority=BelowNormal affinity=0xFF00 peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) first_present=$($PresentLine.Line) exit=$ExitCode"
+        if($ExitCode-ne0){throw "Cold run returned $ExitCode`: $($Run.Name)"}
+        break
+    }
+}
@@ -0,0 +1,57 @@
+[CmdletBinding()]
+param([string]$SourceExecutable='VoxelForgeGpuFieldStep8SkirtViewerB.exe')
+
+$ErrorActionPreference='Stop'
+$HostProcess=Get-Process -Id $PID
+$HostProcess.PriorityClass='BelowNormal'
+$HostProcess.ProcessorAffinity=[IntPtr]0xFF00
+$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..'))
+$BinRoot=Join-Path $StandaloneRoot 'bin'
+$SourcePath=Join-Path $BinRoot $SourceExecutable
+if(-not(Test-Path -LiteralPath $SourcePath -PathType Leaf)){throw "Source viewer not found: $SourcePath"}
+$Runs=@(
+    @{Name='step8_visible_first_blocking';WindowStyle='Normal'},
+    @{Name='step8_hidden_first_blocking';WindowStyle='Hidden'}
+)
+foreach($Run in $Runs)
+{
+    $ExecutableName=if($Run.WindowStyle-eq'Normal')
+        {'VoxelForgeGpuFieldStep8VisibleFirst.exe'}else{'VoxelForgeGpuFieldStep8HiddenFirst.exe'}
+    $ExecutablePath=Join-Path $BinRoot $ExecutableName
+    Copy-Item -LiteralPath $SourcePath -Destination $ExecutablePath -Force
+    $Arguments=@('--gpu-far','--blocking-startup-present','--clip-radius=3','--workers=8','--quit-after-clipmap')
+    $StdoutPath=Join-Path $PSScriptRoot ($Run.Name+'.stdout.log')
+    $StderrPath=Join-Path $PSScriptRoot ($Run.Name+'.stderr.log')
+    Write-Output "VISIBILITY_CONTROL command=`"$ExecutablePath`" $($Arguments -join ' ') window_style=$($Run.WindowStyle)"
+    for($Attempt=0;$Attempt-lt2;$Attempt++)
+    {
+        if($Attempt-gt0)
+        {Write-Output '0xC0000142 encountered; waiting 60 seconds before the single permitted retry.';Start-Sleep -Seconds 60}
+        Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue
+        $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot `
+            -WindowStyle $Run.WindowStyle -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath
+        $PeakPrivateMemory=0L;$OverMemoryLimit=$false
+        while(-not$Child.HasExited)
+        {
+            $Child.Refresh();if($Child.HasExited){break}
+            $Child.PriorityClass='BelowNormal';$Child.ProcessorAffinity=[IntPtr]0xFF00
+            $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()
+        $ExitCode=0
+        if($null-ne$Child.ExitCode){$ExitCode=[int]$Child.ExitCode}
+        if($ExitCode-eq-1073741502-and$Attempt-eq0){continue}
+        if($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=$ExitCode"}
+        $PresentLine=Select-String -Path $StdoutPath -Pattern '^\[PresentProfile\] ordinal=1 '
+        Write-Output "VISIBILITY_EXIT name=$($Run.Name) window_style=$($Run.WindowStyle) pid=$($Child.Id) priority=BelowNormal affinity=0xFF00 peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) first_present=$($PresentLine.Line) exit=$ExitCode"
+        if($ExitCode-ne0){throw "Run returned $ExitCode`: $($Run.Name)"}
+        break
+    }
+}
@@ -0,0 +1,16 @@
+# Step 8 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.6077 | 1.8665 | 3 | 7 | 123 | 42.64% | 0.60% | 0.02% |
+| tunnel | 0.6383 | 3.3799 | 3 | 8 | 130 | 31.73% | 0.87% | 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.9309 | 2.4737 | 5 | 13 | 100 | 43.02% | 2.53% | 0.05% |
+| tunnel_bend | 0.8129 | 2.1599 | 5 | 11 | 55 | 40.96% | 1.80% | 0.07% |
+| remote_yaw0 | 4.2933 | 15.9282 | 61 | 91 | 125 | 15.59% | 9.78% | 6.42% |
+| remote_yaw90 | 1.5692 | 8.8960 | 2 | 68 | 96 | 6.89% | 4.07% | 2.45% |
+| remote_yaw180 | 1.6511 | 8.8623 | 8 | 69 | 102 | 10.33% | 4.90% | 2.69% |
+| remote_yaw270 | 1.7465 | 9.3186 | 3 | 68 | 107 | 10.86% | 4.18% | 2.97% |
@@ -0,0 +1,16 @@
+# step8_clamped 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.6077 | 1.8665 | 3 | 7 | 123 | 42.64% | 0.60% | 0.02% |
+| tunnel | 0.6383 | 3.3799 | 3 | 8 | 130 | 31.73% | 0.87% | 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.9309 | 2.4737 | 5 | 13 | 100 | 43.02% | 2.53% | 0.05% |
+| tunnel_bend | 0.8129 | 2.1599 | 5 | 11 | 55 | 40.96% | 1.80% | 0.07% |
+| remote_yaw0 | 4.2933 | 15.9282 | 61 | 91 | 125 | 15.59% | 9.78% | 6.42% |
+| remote_yaw90 | 1.5692 | 8.8960 | 2 | 68 | 96 | 6.89% | 4.07% | 2.45% |
+| remote_yaw180 | 1.6511 | 8.8623 | 8 | 69 | 102 | 10.33% | 4.90% | 2.69% |
+| remote_yaw270 | 1.7465 | 9.3186 | 3 | 68 | 107 | 10.86% | 4.18% | 2.97% |
@@ -0,0 +1,13 @@
+# Step 8 remote yaw 0, isolated LOD render comparison
+
+The viewer generated and fully-resident clip radius 3 for each pair; `--capture-only-lod` filtered drawing only.
+
+| LOD | MAE RGB | RMSE RGB | Changed pixels | Pixels >32 | Max channel |
+|---:|---:|---:|---:|---:|---:|
+| 2 | 1.3706 | 9.6463 | 3.89% | 1.84% | 125 |
+| 3 | 6.6178 | 20.8273 | 15.13% | 9.96% | 124 |
+| 4 | 5.0988 | 17.4692 | 11.79% | 8.38% | 87 |
+| 5 | 2.6833 | 11.7094 | 8.25% | 4.49% | 87 |
+| 6 | 2.8046 | 11.0939 | 8.97% | 5.66% | 69 |
+| 7 | 1.0024 | 4.3843 | 10.58% | 0.00% | 28 |
+| 8 | 0.2672 | 2.1540 | 1.64% | 0.00% | 21 |
@@ -60,7 +60,9 @@ cbuffer FarConnectorConstants : register(b2)
    uint FarConnectorVertexCount;
    uint FarConnectorPanelVertexCapacity;
    uint FarShapeCount;
-    uint3 FarConnectorPadding;
+    uint FarSkirtFaceMask;
+    float FarSkirtDepthVoxels;
+    uint FarConnectorPadding;
 };

 static const int3 CornerOffset[8] = {
@@ -230,6 +232,8 @@ RWStructuredBuffer<FarConnectorRing> FarConnectorRingOutput : register(u7);
 RWByteAddressBuffer FarConnectorDrawArgs : register(u4);
 RWByteAddressBuffer FarConnectorVertexBytes : register(u5);
 RWByteAddressBuffer FarConnectorIndexBytes : register(u6);
+RWStructuredBuffer<uint> FarSkirtScratch : register(u0);
+ByteAddressBuffer FarSurfaceVertexInput : register(t19);

 RWStructuredBuffer<uint> FarFieldEvalCounters : register(u6);

@@ -510,6 +514,7 @@ float FarRockDistanceAt(float3 position,int3 cachedBinCoord,FarFieldSpatialBin c
 {
    float cave = 100000.0f;
    float overlay = -100000.0f;
+    bool criticalAir=false;
    const int3 positionBinCoord=int3(floor(position/max(1.0f,(float)FarSpatialBinSize)));
    FarFieldSpatialBin range;
    if(all(positionBinCoord==cachedBinCoord))range=cachedRange;
@@ -520,6 +525,25 @@ float FarRockDistanceAt(float3 position,int3 cachedBinCoord,FarFieldSpatialBin c
        const VF_FieldPrimitive shape = FarShapes[FarSpatialBinShapeIds[range.offset + item]];
        if ((shape.flags & VF_FIELD_PRIMITIVE_CARVE) != 0)
        {
+            // The CPU field forces critical passage air at arbitrary skirt
+            // probes using the containing integer voxel. Match that decided
+            // exception here; the ordinary smooth field remains GPU-evaluated.
+            if((shape.flags&VF_FIELD_PRIMITIVE_CRITICAL)!=0
+                &&all(position>=float3(shape.minVoxel[0],shape.minVoxel[1],shape.minVoxel[2]))
+                &&all(position<=float3(shape.maxVoxel[0],shape.maxVoxel[1],shape.maxVoxel[2])))
+            {
+                const int3 decidedPosition=int3(position);
+                float decidedDistance=VF_FieldPrimitiveDistance(shape,
+                    (float)decidedPosition.x,(float)decidedPosition.y,(float)decidedPosition.z);
+                if((shape.flags&VF_FIELD_PRIMITIVE_HAS_FLOOR)!=0)
+                {
+                    const float decidedFloor=VF_FieldPrimitiveFloorAt(shape,
+                        (float)decidedPosition.x,(float)decidedPosition.y);
+                    decidedDistance=VF_FieldSmoothMax(decidedDistance,
+                        decidedFloor-(float)decidedPosition.z,shape.floorBlendRadius);
+                }
+                criticalAir=criticalAir||decidedDistance<=0.0f;
+            }
            if (position.x < shape.minVoxel[0] - 6 || position.x > shape.maxVoxel[0] + 6
                || position.y < shape.minVoxel[1] - 6 || position.y > shape.maxVoxel[1] + 6
                || position.z < shape.minVoxel[2] - 6 || position.z > shape.maxVoxel[2] + 6)
@@ -541,7 +565,174 @@ float FarRockDistanceAt(float3 position,int3 cachedBinCoord,FarFieldSpatialBin c
        }
    }
    const float base = min(position.z - FarStrataLo, FarStrataHi - position.z);
-    return VF_FieldSmoothMin(base, max(cave, overlay), FarStrataBlend);
+    float rockDistance=VF_FieldSmoothMin(base,max(cave,overlay),FarStrataBlend);
+    if(criticalAir)rockDistance=min(rockDistance,-1.0f/16.0f);
+    return rockDistance;
+}
+
+struct FarSkirtVertex
+{
+    float3 position;
+    float3 normal;
+    float density;
+};
+
+float FarSkirtDensityAtVoxel(float3 position)
+{
+    const int3 binCoord=int3(floor(position/max(1.0f,(float)FarSpatialBinSize)));
+    const FarFieldSpatialBin range=FarFieldFindBin(position);
+    const float rockDistance=FarRockDistanceAt(position,binCoord,range);
+    float density=-rockDistance;
+    // ShapeWorld clamps arbitrary (off-lattice) density magnitudes near the
+    // surface to 0.05 voxel. Skirt clipping uses those values for edge-root
+    // interpolation, so matching only their signs can move a clipped edge.
+    if(abs(density)<0.05f)density=rockDistance>=0.0f?-0.05f:0.05f;
+    return density;
+}
+
+float FarSkirtDensityAtCm(float3 positionCm)
+{
+    return FarSkirtDensityAtVoxel(positionCm/VoxelSizeCm);
+}
+
+FarSkirtVertex FarSkirtInterpolate(FarSkirtVertex a,FarSkirtVertex b,float t)
+{
+    if(t<=1.0e-6f)return a;
+    if(t>=1.0f-1.0e-6f)return b;
+    FarSkirtVertex result;
+    result.position=lerp(a.position,b.position,t);
+    result.normal=normalize(lerp(a.normal,b.normal,t));
+    result.density=FarSkirtDensityAtCm(result.position);
+    return result;
+}
+
+void FarSkirtClipTriangle(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex c,
+    out FarSkirtVertex polygon[5],out uint polygonCount)
+{
+    FarSkirtVertex input[3];input[0]=a;input[1]=b;input[2]=c;
+    polygonCount=0;
+    [unroll]
+    for(uint edge=0;edge<3;++edge)
+    {
+        FarSkirtVertex start=input[edge];
+        FarSkirtVertex end=input[(edge+1)%3];
+        const bool startSolid=start.density<=0.0f;
+        const bool endSolid=end.density<=0.0f;
+        if(startSolid!=endSolid)
+        {
+            const float denominator=start.density-end.density;
+            const float t=abs(denominator)>1.0e-8f
+                ?clamp(start.density/denominator,0.0f,1.0f):0.5f;
+            polygon[polygonCount++]=FarSkirtInterpolate(start,end,t);
+        }
+        if(endSolid)polygon[polygonCount++]=end;
+    }
+}
+
+uint FarSkirtSourceTriangleOutputCount(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex c)
+{
+    FarSkirtVertex polygon[5];uint polygonCount=0;
+    FarSkirtClipTriangle(a,b,c,polygon,polygonCount);
+    if(polygonCount<3)return 0;
+    uint vertexCount=0;
+    [loop]
+    for(uint fan=1;fan+1<polygonCount;++fan)
+    {
+        const float3 centerVoxel=(polygon[0].position+polygon[fan].position
+            +polygon[fan+1].position)/(3.0f*VoxelSizeCm);
+        if(FarSkirtDensityAtVoxel(centerVoxel)<=0.0f)vertexCount+=6;
+    }
+    return vertexCount;
+}
+
+uint FarSkirtBoundaryFaceMask(float3 a,float3 b)
+{
+    const float3 tileMin=float3(OriginX,OriginY,OriginZ)*VoxelSizeCm;
+    const float3 tileMax=(float3(OriginX,OriginY,OriginZ)
+        +float(CellsPerAxis*StepVoxels))*VoxelSizeCm;
+    uint mask=0;
+    if(a.x==tileMin.x&&b.x==tileMin.x)mask|=1u<<0;
+    if(a.x==tileMax.x&&b.x==tileMax.x)mask|=1u<<1;
+    if(a.y==tileMin.y&&b.y==tileMin.y)mask|=1u<<2;
+    if(a.y==tileMax.y&&b.y==tileMax.y)mask|=1u<<3;
+    if(a.z==tileMin.z&&b.z==tileMin.z)mask|=1u<<4;
+    if(a.z==tileMax.z&&b.z==tileMax.z)mask|=1u<<5;
+    return mask;
+}
+
+FarSkirtVertex FarSkirtSurfaceVertex(uint vertexId)
+{
+    const uint byteAddress=(OutputBaseVertex+vertexId)*32;
+    FarSkirtVertex vertex;
+    vertex.position=asfloat(FarSurfaceVertexInput.Load3(byteAddress));
+    vertex.normal=asfloat(FarSurfaceVertexInput.Load3(byteAddress+12));
+    vertex.density=0.0f;
+    return vertex;
+}
+
+FarSkirtVertex FarSkirtExtrudedVertex(FarSkirtVertex source)
+{
+    FarSkirtVertex vertex=source;
+    vertex.position-=source.normal*(max(1.0f,FarSkirtDepthVoxels)*VoxelSizeCm);
+    vertex.density=FarSkirtDensityAtCm(vertex.position);
+    return vertex;
+}
+
+uint FarSkirtEdgeOutputCount(uint edgeVertexId)
+{
+    const uint surfaceVertexCount=min(FarConnectorDrawArgs.Load(OutputArgsByteOffset),OutputVertexCapacity);
+    if(edgeVertexId>=surfaceVertexCount)return 0;
+    const uint triangleBase=edgeVertexId-(edgeVertexId%3);
+    const uint edge=edgeVertexId%3;
+    const FarSkirtVertex a=FarSkirtSurfaceVertex(triangleBase+edge);
+    const FarSkirtVertex b=FarSkirtSurfaceVertex(triangleBase+(edge+1)%3);
+    if((FarSkirtBoundaryFaceMask(a.position,b.position)&FarSkirtFaceMask)==0)return 0;
+    const FarSkirtVertex a2=FarSkirtExtrudedVertex(a);
+    const FarSkirtVertex b2=FarSkirtExtrudedVertex(b);
+    return FarSkirtSourceTriangleOutputCount(a,b,b2)
+        +FarSkirtSourceTriangleOutputCount(a,b2,a2);
+}
+
+void FarSkirtStoreVertex(uint outputId,FarSkirtVertex vertex,uint panelBaseVertex)
+{
+    GpuVertex output;
+    output.position=vertex.position;
+    output.normal=vertex.normal;
+    output.uv=vertex.position.xy/VoxelSizeCm;
+    const uint index=panelBaseVertex+outputId;
+    const uint byteAddress=index*32;
+    FarConnectorVertexBytes.Store3(byteAddress,asuint(output.position));
+    FarConnectorVertexBytes.Store3(byteAddress+12,asuint(output.normal));
+    FarConnectorVertexBytes.Store2(byteAddress+24,asuint(output.uv));
+    FarConnectorIndexBytes.Store(index*4,index);
+}
+
+void FarSkirtEmitTriangle(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex c,
+    uint panelBaseVertex,uint edgeBase,inout uint edgeCursor)
+{
+    FarSkirtStoreVertex(edgeBase+edgeCursor++,a,panelBaseVertex);
+    FarSkirtStoreVertex(edgeBase+edgeCursor++,b,panelBaseVertex);
+    FarSkirtStoreVertex(edgeBase+edgeCursor++,c,panelBaseVertex);
+    FarSkirtStoreVertex(edgeBase+edgeCursor++,a,panelBaseVertex);
+    FarSkirtStoreVertex(edgeBase+edgeCursor++,c,panelBaseVertex);
+    FarSkirtStoreVertex(edgeBase+edgeCursor++,b,panelBaseVertex);
+}
+
+void FarSkirtEmitSourceTriangle(FarSkirtVertex a,FarSkirtVertex b,FarSkirtVertex c,
+    uint panelBaseVertex,uint edgeBase,inout uint edgeCursor)
+{
+    FarSkirtVertex polygon[5];uint polygonCount=0;
+    FarSkirtClipTriangle(a,b,c,polygon,polygonCount);
+    if(polygonCount<3)return;
+    [loop]
+    for(uint fan=1;fan+1<polygonCount;++fan)
+    {
+        const float3 centerVoxel=(polygon[0].position+polygon[fan].position
+            +polygon[fan+1].position)/(3.0f*VoxelSizeCm);
+        if(FarSkirtDensityAtVoxel(centerVoxel)<=0.0f)
+            FarSkirtEmitTriangle(polygon[0],polygon[fan],polygon[fan+1],
+                panelBaseVertex,edgeBase,edgeCursor);
+    }
 }

 float FarConnectorRadius(float3 center, float3 direction, float nominalRadius)
@@ -704,6 +895,7 @@ void BuildFarConnectorRingsCS(uint3 groupId : SV_GroupID,
            FarConnectorDrawArgs.Store(OutputArgsByteOffset + 28, panelBaseVertex);
            FarConnectorDrawArgs.Store(OutputArgsByteOffset + 32, 0);
            FarConnectorDrawArgs.Store(OutputArgsByteOffset + 36, 0);
+            FarConnectorDrawArgs.Store(OutputArgsByteOffset + 40, 0);
        }
    }
    GroupMemoryBarrierWithGroupSync();
@@ -880,6 +1072,7 @@ void ScanGroupsCS(uint3 dispatchId : SV_DispatchThreadID)
    DrawArgs.Store(OutputArgsByteOffset + 28, 0);
    DrawArgs.Store(OutputArgsByteOffset + 32, 0);
    DrawArgs.Store(OutputArgsByteOffset + 36, 0);
+    DrawArgs.Store(OutputArgsByteOffset + 40, 0);
 }

 StructuredBuffer<float> Density : register(t2);
@@ -927,6 +1120,63 @@ void EmitCS(uint3 dispatchId : SV_DispatchThreadID)
    }
 }

+[numthreads(64,1,1)]
+void ClassifyFarSkirtsCS(uint3 dispatchId : SV_DispatchThreadID)
+{
+    const uint edgeVertexId=dispatchId.x;
+    if(edgeVertexId>=OutputVertexCapacity)return;
+    FarSkirtScratch[edgeVertexId]=FarSkirtEdgeOutputCount(edgeVertexId);
+}
+
+[numthreads(1,1,1)]
+void ScanFarSkirtsCS(uint3 dispatchId : SV_DispatchThreadID)
+{
+    const uint surfaceVertexCount=min(FarConnectorDrawArgs.Load(OutputArgsByteOffset),
+        OutputVertexCapacity);
+    uint totalSkirtVertices=0;
+    [loop]
+    for(uint edgeVertexId=0;edgeVertexId<surfaceVertexCount;++edgeVertexId)
+    {
+        const uint count=FarSkirtScratch[edgeVertexId];
+        FarSkirtScratch[edgeVertexId]=totalSkirtVertices;
+        totalSkirtVertices+=count;
+    }
+    const uint panelVertexCount=FarConnectorDrawArgs.Load(OutputArgsByteOffset+20);
+    if(panelVertexCount+totalSkirtVertices>FarConnectorPanelVertexCapacity)
+    {
+        FarConnectorDrawArgs.Store(OutputArgsByteOffset+40,1);
+        return;
+    }
+    FarConnectorDrawArgs.Store(OutputArgsByteOffset+40,0);
+    FarConnectorDrawArgs.Store(OutputArgsByteOffset+20,panelVertexCount+totalSkirtVertices);
+    FarConnectorDrawArgs.Store(OutputArgsByteOffset+24,
+        panelVertexCount+totalSkirtVertices>0?1:0);
+}
+
+[numthreads(64,1,1)]
+void EmitFarSkirtsCS(uint3 dispatchId : SV_DispatchThreadID)
+{
+    const uint edgeVertexId=dispatchId.x;
+    if(edgeVertexId>=OutputVertexCapacity||FarSkirtFaceMask==0
+        ||FarConnectorDrawArgs.Load(OutputArgsByteOffset+40)!=0)return;
+    const uint surfaceVertexCount=min(FarConnectorDrawArgs.Load(OutputArgsByteOffset),
+        OutputVertexCapacity);
+    if(edgeVertexId>=surfaceVertexCount)return;
+    if(FarSkirtEdgeOutputCount(edgeVertexId)==0)return;
+    const uint triangleBase=edgeVertexId-(edgeVertexId%3);
+    const uint edge=edgeVertexId%3;
+    const FarSkirtVertex a=FarSkirtSurfaceVertex(triangleBase+edge);
+    const FarSkirtVertex b=FarSkirtSurfaceVertex(triangleBase+(edge+1)%3);
+    if((FarSkirtBoundaryFaceMask(a.position,b.position)&FarSkirtFaceMask)==0)return;
+    const FarSkirtVertex a2=FarSkirtExtrudedVertex(a);
+    const FarSkirtVertex b2=FarSkirtExtrudedVertex(b);
+    const uint panelBaseVertex=FarConnectorDrawArgs.Load(OutputArgsByteOffset+28);
+    const uint edgeBase=FarConnectorVertexCount+FarSkirtScratch[edgeVertexId];
+    uint edgeCursor=0;
+    FarSkirtEmitSourceTriangle(a,b,b2,panelBaseVertex,edgeBase,edgeCursor);
+    FarSkirtEmitSourceTriangle(a,b2,a2,panelBaseVertex,edgeBase,edgeCursor);
+}
+
 // Smoke probe for the dual-language noise path. This is intentionally separate from the mesh
 // dispatch until shape records can be exported from the exact CPU layout compiler.
 struct NoiseProbeInput { float3 position; uint seed; };
@@ -1361,6 +1361,8 @@ namespace
        VoxelForgeGpuFarViewer::GpuTimingQueries TimingQueries{};
        uint64 VertexCapacityBytes=0;
        uint32 OutputSlot=0,IndirectArgsByteOffset=0,BatchIndex=UINT32_MAX;
+        uint8 VisibleSkirtFaceMask=0;
+        bool bSkirtDrawArgsAudited=false;
        double GpuDispatchMs=0.0;
        bool bGpuComplete=false;
        bool bTimingResolved=false;
@@ -1486,13 +1488,26 @@ namespace
        bool bGpuFarEnabled=false;
        bool bGpuFarAuditEnabled=false;
        bool bGpuFarAuditReported=false;
+        bool bGpuSkirtEmitEnabled=true,bGpuConnectorRepairEnabled=true;
        bool bDisableVSync=false;
+        bool bForceBlockingStartupPresent=false;
+        bool bPresentCpuFrameBeforeGpu=false;
+        bool bCpuReadyFramePresentedBeforeGpu=false;
+        bool bHoldCpuFrameDuringGpuWarmup=false,bGpuWarmupHoldReported=false;
+        bool bGpuWarmupReleaseReported=false;
+        uint64 GpuWarmupHeldPresents=0;
+        bool bAuditGpuSkirtDrawArgs=false,bGpuSkirtDrawAuditReported=false;
+        ID3D11Buffer* GpuSkirtArgsStaging=nullptr;
+        uint64 GpuSkirtDrawAuditTiles=0,GpuSkirtDrawAuditZeroCounts=0,
+            GpuSkirtDrawAuditOverflow=0,GpuSkirtDrawAuditBadStartIndex=0;
+        uint64 GpuSkirtDrawAuditSumIndices=0,GpuSkirtDrawAuditMaxIndices=0;
        std::unique_ptr<FFarSurvivalAudit> GpuFarAudit;
        bool bGpuFarFirstTileReported=false;
        std::filesystem::path GpuFarShaderPath;
        double GpuFarEnqueueWallSeconds=0.0;
        uint64 GpuFarTilesSubmitted=0;
        uint64 PresentCallCount=0;
+        uint64 StartupNonBlockingPresentCalls=0,StartupNonBlockingPresentSkips=0;
        double GpuFarGpuDispatchMs=0.0;
        double GpuFarGpuDispatchMaxMs=0.0;
        double GpuFarSubmissionSpanMs=0.0,GpuFarSubmissionSpanMaxMs=0.0;
@@ -1549,6 +1564,10 @@ namespace
        int32 ClipRadius=3;
        int32 MaxClipLevelOverride=-1;
        int32 CaptureTileLimit=0;
+        int32 CaptureOnlyLod=-1;
+        bool bCaptureOnlyTile=false;
+        int32 CaptureOnlyTileX=0,CaptureOnlyTileY=0,CaptureOnlyTileZ=0,CaptureOnlyTileLod=0;
+        int32 CaptureOutputSlot=-1;
        int32 Width=1280,Height=800;
        float Yaw=0.0f,Pitch=0.0f,MoveSpeed=8.0f;
        V3 Position{};
@@ -1565,6 +1584,7 @@ namespace
            StopGeneration(true);ReleaseMeshes();
            if(SolidRasterizer)SolidRasterizer->Release();if(WireRasterizer)WireRasterizer->Release();
            if(Constants)Constants->Release();if(InputLayout)InputLayout->Release();
+            if(GpuSkirtArgsStaging)GpuSkirtArgsStaging->Release();
            if(VertexShader)VertexShader->Release();if(PixelShader)PixelShader->Release();
            if(RenderTarget)RenderTarget->Release();if(DepthView)DepthView->Release();
            if(SwapChain)SwapChain->Release();if(Context)Context->Release();if(Device)Device->Release();
@@ -2101,6 +2121,24 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                for(const auto& Key:CoreTiles)
                    DesiredTiles.emplace_back(FIntVector(Key.coordinate.x,Key.coordinate.y,Key.coordinate.z),Key.lod);
            }
+            if(bCaptureOnlyTile)
+            {
+                const size_t Before=DesiredTiles.size();
+                DesiredTiles.erase(std::remove_if(DesiredTiles.begin(),DesiredTiles.end(),[&](const FVoxelTileKey& Key)
+                {
+                    return Key.Level!=CaptureOnlyTileLod||Key.Coord.X!=CaptureOnlyTileX
+                        ||Key.Coord.Y!=CaptureOnlyTileY||Key.Coord.Z!=CaptureOnlyTileZ;
+                }),DesiredTiles.end());
+                std::printf("[CaptureOnlyTile] key=(%d,%d,%d,lod%d) before=%zu selected=%zu\n",
+                    CaptureOnlyTileX,CaptureOnlyTileY,CaptureOnlyTileZ,CaptureOnlyTileLod,
+                    Before,DesiredTiles.size());
+                if(DesiredTiles.empty())
+                {
+                    Error="the requested capture-only tile is not in the desired clipmap";
+                    std::fprintf(stderr,"World load failed: %s\n",Error.c_str());
+                    return;
+                }
+            }
            if(CaptureTileLimit>0&&CapturePath.size())
            {
                const size_t Before=DesiredTiles.size();
@@ -2155,7 +2193,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                }
            }
            if(bShapePipelineActive)Streamer->SetDesiredTileFilter(InitialNeighborhoodTiles);
-            else if(CaptureTileLimit>0&&CapturePath.size())
+            else if((CaptureTileLimit>0&&CapturePath.size())||bCaptureOnlyTile)
                Streamer->SetDesiredTileFilter(DesiredTiles);
            const FStandaloneStreamingUpdate Update=Streamer->Start(CameraVoxels,{},CrossingIndex);
            if(bShapePipelineActive&&Update.DesiredTiles!=static_cast<int32>(InitialNeighborhoodTiles.size()))
@@ -2343,6 +2381,20 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                            Inputs[Index]=Key.Level<World->FullResClipLevels
                                ?ShapeWorld->exportGpuFieldTile(CoreKey,Errors[Index])
                                :ShapeWorld->exportFarFieldTile(CoreKey,Errors[Index]);
+                            if(Inputs[Index]&&Key.Level>=World->FullResClipLevels)
+                            {
+                                Inputs[Index]->visibleSkirtFaceMask=World->bGenerateSkirts
+                                    &&bGpuSkirtEmitEnabled
+                                    ?ComputeVisibleSkirtFaceMask(Key,DesiredTileSet):0;
+                                Inputs[Index]->skirtDepthVoxels=std::max(1.0f,
+                                    static_cast<float>(World->SkirtCells))*Key.StepVoxels();
+                                if(!bGpuConnectorRepairEnabled)
+                                {
+                                    Inputs[Index]->connectorPanels.clear();
+                                    Inputs[Index]->connectorRingCount=0;
+                                    Inputs[Index]->connectorVertexCount=0;
+                                }
+                            }
                        }
                        catch(const std::exception& Error) { Errors[Index]=Error.what(); }
                    }
@@ -2390,12 +2442,22 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                {
                    const uint8 VisibleSkirtFaces=World->bGenerateSkirts
                        ?ComputeVisibleSkirtFaceMask(Key,DesiredTileSet):0;
-                    if(VisibleSkirtFaces!=0)
+                    if(Key.Level>=World->FullResClipLevels)
                    {
-                        // CPU tiles carry the established clipped LOD skirts. The GPU field
-                        // and marching-cubes outputs match the CPU surface, but currently do
-                        // not emit these seam curtains. Keep tiles with exposed LOD faces on
-                        // the reference mesh path until a GPU skirt kernel is available.
+                        const auto Existing=GpuFarTiles.find(Key);
+                        if(Existing!=GpuFarTiles.end()
+                            &&Existing->second.VisibleSkirtFaceMask!=VisibleSkirtFaces)
+                        {
+                            const uint32 ReusedSlot=Existing->second.OutputSlot;
+                            ReleaseGpuFarTile(Existing->second);
+                            GpuFarTiles.erase(Existing);
+                            GpuFarFreeOutputSlots.push_back(ReusedSlot);
+                        }
+                    }
+                    if(VisibleSkirtFaces!=0&&NearGpuEligible)
+                    {
+                        // Keep fine tiles on the CPU reference path. Far tiles carry the
+                        // CPU-decided visible face mask to the deterministic GPU skirt pass.
                        CpuTiles.push_back(Key);
                        ++SkirtTransitionFallbackTiles;
                    }
@@ -2411,6 +2473,14 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
            }
            if(!Eligible.empty()&&GpuFarOutputCapacitySlots==0)
                GpuFarOutputCapacitySlots=static_cast<uint32>(Eligible.size());
+            if(bCaptureOnlyTile&&CaptureOutputSlot>=0)
+            {
+                GpuFarOutputCapacitySlots=std::max(GpuFarOutputCapacitySlots,
+                    static_cast<uint32>(CaptureOutputSlot+1));
+                GpuFarNextOutputSlot=static_cast<uint32>(CaptureOutputSlot);
+                std::printf("[CaptureOutputSlotProbe] slot=%d pool_slots=%u\n",
+                    CaptureOutputSlot,GpuFarOutputCapacitySlots);
+            }
            if(!Eligible.empty())
            {
                std::string PoolError;const auto ReserveStart=std::chrono::steady_clock::now();
@@ -2473,7 +2543,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                    {return Key.Level<World->FullResClipLevels;}));
            std::printf("[GpuFieldEligibility] eligible_gpu_tiles=%zu near_gpu_tiles=%zu cpu_fallback_tiles=%zu\n",
                GpuFarEligibleTiles.size(),NearGpuTiles,Keys.size()-GpuFarEligibleTiles.size());
-            std::printf("[GpuSkirtSeamFallback] tiles=%llu reason=visible_lod_skirt_faces status=active\n",
+            std::printf("[GpuSkirtSeamFallback] near_tiles=%llu reason=fine_visible_lod_skirts gpu_far_skirts=active\n",
                static_cast<unsigned long long>(SkirtTransitionFallbackTiles));
        }

@@ -2533,6 +2603,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                    bGpuFarEnabled=false;
                    break;
                }
+                const uint8 PreparedSkirtFaceMask=PreparedInput->visibleSkirtFaceMask;
                if(!VoxelForgeGpuFarViewer::BuildFarTile(Device,Context,*ShapeWorld,CoreKey,
                    PreparedInput,
                    GpuFarShaderPath,&Tile.VertexBuffer,&Tile.IndexBuffer,&Tile.IndirectArgs,
@@ -2565,6 +2636,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                }
                Tile.VertexCapacityBytes=static_cast<uint64>(
                    VoxelForgeGpuFarViewer::PooledOutputVertexCapacityPerTile)*32u;
+                Tile.VisibleSkirtFaceMask=PreparedSkirtFaceMask;
                Tile.GpuDispatchMs=GpuDispatchMs;
                GpuFarEnqueueWallSeconds+=EnqueueMs/1000.0;
                if(GpuDispatchMs>0.0)
@@ -2643,6 +2715,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
            bGpuFarEnabled=!bGpuFarEnabled;
            if(bGpuFarEnabled)
            {
+                bPresentCpuFrameBeforeGpu=true;
+                bHoldCpuFrameDuringGpuWarmup=true;
+                bCpuReadyFramePresentedBeforeGpu=false;
                ScheduleGpuFarTiles(DesiredTiles);
                std::vector<FVoxelTileKey> CpuFar;
                for(const FVoxelTileKey& Key:DesiredTiles)
@@ -3241,6 +3316,10 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                    static_cast<unsigned long long>(TotalIndices),StreamingBatches.size());
                std::printf("[ColdTimeline] event=whole_clipmap_ready at_ms=%.3f waits_for=%zu_GPU_tiles_plus_CPU_fallback_and_uploads\n",
                    Elapsed*1000.0,GpuFarEligibleTiles.size());
+                std::printf("[PresentStartupSummary] nonblocking_startup_calls=%llu skipped_busy_calls=%llu gpu_warmup_held_calls=%llu policy=do_not_wait_until_initial_clipmap_ready\n",
+                    static_cast<unsigned long long>(StartupNonBlockingPresentCalls),
+                    static_cast<unsigned long long>(StartupNonBlockingPresentSkips),
+                    static_cast<unsigned long long>(GpuWarmupHeldPresents));
                if(bGpuFarEnabled)
                {
                    SampleGpuFarMemory();
@@ -3307,7 +3386,10 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;

        void UpdateTitle()
        {
-            uint64 VisibleIndices=0;for(const auto& Pair:StreamingBatches)if(bShowLevel[Pair.first.Level])VisibleIndices+=Pair.second.IndexCount;
+            uint64 VisibleIndices=0;for(const auto& Pair:StreamingBatches)
+                if(bShowLevel[Pair.first.Level]
+                    &&(CaptureOnlyLod<0||Pair.first.Level==CaptureOnlyLod))
+                    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 | field "<<(bGpuFarEnabled?"GPU":"CPU")
@@ -3323,7 +3405,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                ?std::chrono::duration<double,std::milli>(DrawStart-LastDrawEndAt).count():0.0;
            PumpGeneration();
            const auto CpuPumpEnd=std::chrono::steady_clock::now();
-            PumpGpuFarGeneration();
+            const bool bDeferGpuUntilCpuFrame=bGpuFarEnabled&&bPresentCpuFrameBeforeGpu
+                &&!bCpuReadyFramePresentedBeforeGpu;
+            if(!bDeferGpuUntilCpuFrame)PumpGpuFarGeneration();
            const auto GpuPumpEnd=std::chrono::steady_clock::now();
            const auto Now=std::chrono::steady_clock::now();
            if(FlightTargetCrossings>0&&bFullDetailReported&&!bFlightStarted)
@@ -3373,7 +3457,9 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
            for(const auto& Pair:StreamingBatches)
            {
                const FStreamingGpuBatch& L=Pair.second;
-                if(!bShowLevel[Pair.first.Level]||!L.VertexBuffer||!L.IndexBuffer||!L.IndexCount)continue;
+                if(!bShowLevel[Pair.first.Level]
+                    ||(CaptureOnlyLod>=0&&Pair.first.Level!=CaptureOnlyLod)
+                    ||!L.VertexBuffer||!L.IndexBuffer||!L.IndexCount)continue;
                const UINT Stride=sizeof(FGpuVertex),Offset=0;Context->IASetVertexBuffers(0,1,&L.VertexBuffer,&Stride,&Offset);
                Context->IASetIndexBuffer(L.IndexBuffer,DXGI_FORMAT_R32_UINT,0);Context->DrawIndexed(L.IndexCount,0,0);
            }
@@ -3385,12 +3471,59 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                C.Padding[0]=0.01f;
                Context->UpdateSubresource(Constants,0,nullptr,&C,0,0);
                Context->VSSetConstantBuffers(0,1,&Constants);Context->PSSetConstantBuffers(0,1,&Constants);
-                for(const auto& Pair:GpuFarTiles)
+                for(auto& Pair:GpuFarTiles)
                {
-                    const FGpuFarTile& Tile=Pair.second;
-                    if(!bShowLevel[Pair.first.Level]||!Tile.VertexBuffer||!Tile.IndexBuffer||!Tile.IndirectArgs)continue;
+                    FGpuFarTile& Tile=Pair.second;
+                    if(!bShowLevel[Pair.first.Level]
+                        ||(CaptureOnlyLod>=0&&Pair.first.Level!=CaptureOnlyLod)
+                        ||!Tile.VertexBuffer||!Tile.IndexBuffer||!Tile.IndirectArgs)continue;
                    if(!bInitialClipmapReported&&Pair.first.Level>=World->FullResClipLevels)continue;
                    if(!IsRenderedTileReady(Pair.first))continue;
+                    if(bAuditGpuSkirtDrawArgs&&Tile.VisibleSkirtFaceMask!=0&&!Tile.bSkirtDrawArgsAudited)
+                    {
+                        if(!GpuSkirtArgsStaging)
+                        {
+                            D3D11_BUFFER_DESC StagingDesc{};
+                            StagingDesc.ByteWidth=11u*sizeof(uint32);
+                            StagingDesc.Usage=D3D11_USAGE_STAGING;
+                            StagingDesc.CPUAccessFlags=D3D11_CPU_ACCESS_READ;
+                            const HRESULT CreateResult=Device->CreateBuffer(&StagingDesc,nullptr,
+                                &GpuSkirtArgsStaging);
+                            if(FAILED(CreateResult))
+                                std::fprintf(stderr,"[GpuSkirtDrawAuditError] create_staging_hr=0x%08lx\n",
+                                    static_cast<unsigned long>(CreateResult));
+                        }
+                        if(GpuSkirtArgsStaging)
+                        {
+                            const D3D11_BOX ArgsSource{Tile.IndirectArgsByteOffset,0,0,
+                                Tile.IndirectArgsByteOffset+11u*sizeof(uint32),1,1};
+                            Context->CopySubresourceRegion(GpuSkirtArgsStaging,0,0,0,0,
+                                Tile.IndirectArgs,0,&ArgsSource);
+                            D3D11_MAPPED_SUBRESOURCE ArgsMapped{};
+                            const HRESULT MapResult=Context->Map(GpuSkirtArgsStaging,0,
+                                D3D11_MAP_READ,0,&ArgsMapped);
+                            if(SUCCEEDED(MapResult))
+                            {
+                                const auto* Args=static_cast<const uint32*>(ArgsMapped.pData);
+                                const uint32 ExpectedStartIndex=(Tile.IndirectArgsByteOffset
+                                    /(11u*sizeof(uint32)))*32768u;
+                                ++GpuSkirtDrawAuditTiles;
+                                GpuSkirtDrawAuditSumIndices+=Args[5];
+                                GpuSkirtDrawAuditMaxIndices=std::max<uint64>(GpuSkirtDrawAuditMaxIndices,Args[5]);
+                                if(Args[5]==0)++GpuSkirtDrawAuditZeroCounts;
+                                if(Args[10]!=0)++GpuSkirtDrawAuditOverflow;
+                                if(Args[7]!=ExpectedStartIndex)++GpuSkirtDrawAuditBadStartIndex;
+                                std::printf("[GpuSkirtDrawArgs] key=(%d,%d,%d,lod%d) face_mask=0x%02x output_slot=%u surface_indices=%u skirt_panel_indices=%u instances=%u connector_start_index=%u expected_start_index=%u overflow=%u\n",
+                                    Pair.first.Coord.X,Pair.first.Coord.Y,Pair.first.Coord.Z,Pair.first.Level,
+                                    static_cast<unsigned>(Tile.VisibleSkirtFaceMask),Tile.OutputSlot,
+                                    Args[0],Args[5],Args[6],Args[7],ExpectedStartIndex,Args[10]);
+                                Context->Unmap(GpuSkirtArgsStaging,0);
+                            }
+                            else std::fprintf(stderr,"[GpuSkirtDrawAuditError] map_staging_hr=0x%08lx\n",
+                                static_cast<unsigned long>(MapResult));
+                        }
+                        Tile.bSkirtDrawArgsAudited=true;
+                    }
                    const UINT Stride=32,Offset=0;
                    Context->IASetVertexBuffers(0,1,&Tile.VertexBuffer,&Stride,&Offset);
                    Context->IASetIndexBuffer(Tile.IndexBuffer,DXGI_FORMAT_R32_UINT,0);
@@ -3404,20 +3537,84 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75;
                    }
                }
            }
+            if(bAuditGpuSkirtDrawArgs&&bInitialClipmapReported&&!bGpuSkirtDrawAuditReported)
+            {
+                std::printf("[GpuSkirtDrawArgsSummary] tiles=%llu zero_index_counts=%llu overflow=%llu bad_slot_start_indices=%llu total_panel_plus_skirt_indices=%llu max_panel_plus_skirt_indices=%llu status=%s\n",
+                    static_cast<unsigned long long>(GpuSkirtDrawAuditTiles),
+                    static_cast<unsigned long long>(GpuSkirtDrawAuditZeroCounts),
+                    static_cast<unsigned long long>(GpuSkirtDrawAuditOverflow),
+                    static_cast<unsigned long long>(GpuSkirtDrawAuditBadStartIndex),
+                    static_cast<unsigned long long>(GpuSkirtDrawAuditSumIndices),
+                    static_cast<unsigned long long>(GpuSkirtDrawAuditMaxIndices),
+                    (GpuSkirtDrawAuditTiles>0&&GpuSkirtDrawAuditOverflow==0
+                        &&GpuSkirtDrawAuditBadStartIndex==0)?"PASS":"FAIL");
+                bGpuSkirtDrawAuditReported=true;
+                std::fflush(stdout);
+            }
            CaptureRenderedFrame();
+            const bool bHoldFrontBufferForGpuWarmup=bHoldCpuFrameDuringGpuWarmup
+                &&bCpuReadyFramePresentedBeforeGpu&&bGpuFarEnabled
+                &&GpuFarTilesSubmitted>0&&!bInitialClipmapReported;
+            if(bHoldFrontBufferForGpuWarmup)
+            {
+                if(!bGpuWarmupHoldReported)
+                {
+                    bGpuWarmupHoldReported=true;
+                    std::printf("[GpuWarmupPresentHold] begin_ms=%.3f policy=retain_last_CPU_ready_front_buffer_until_GPU_clipmap_ready\n",
+                        std::chrono::duration<double,std::milli>(
+                            std::chrono::steady_clock::now()-ColdProcessStart).count());
+                    std::fflush(stdout);
+                }
+                ++GpuWarmupHeldPresents;
+                // Keep pumping CPU and GPU work at the display cadence without
+                // calling the driver present path that stalls on its first GPU frame.
+                std::this_thread::sleep_for(std::chrono::milliseconds(11));
+                LastDrawEndAt=std::chrono::steady_clock::now();
+                return;
+            }
            const auto BeforePresent=std::chrono::steady_clock::now();
-            const UINT PresentSyncInterval=bDisableVSync?0u:1u;
+            const bool bNonBlockingStartupPresent=bGpuFarEnabled&&!bInitialClipmapReported
+                &&!bForceBlockingStartupPresent;
+            const UINT PresentSyncInterval=(bDisableVSync||bNonBlockingStartupPresent)?0u:1u;
+            const UINT PresentFlags=bNonBlockingStartupPresent?DXGI_PRESENT_DO_NOT_WAIT:0u;
            const auto PresentCallStart=std::chrono::steady_clock::now();
-            SwapChain->Present(PresentSyncInterval,0);
+            const HRESULT PresentResult=SwapChain->Present(PresentSyncInterval,PresentFlags);
            const auto PresentCallEnd=std::chrono::steady_clock::now();
+            if(bGpuWarmupHoldReported&&!bGpuWarmupReleaseReported)
+            {
+                bGpuWarmupReleaseReported=true;
+                std::printf("[GpuWarmupPresentRelease] at_ms=%.3f first_post_warmup_present_ms=%.3f sync=%u hr=0x%08lx\n",
+                    std::chrono::duration<double,std::milli>(PresentCallEnd-ColdProcessStart).count(),
+                    std::chrono::duration<double,std::milli>(PresentCallEnd-PresentCallStart).count(),
+                    PresentSyncInterval,static_cast<unsigned long>(PresentResult));
+                std::fflush(stdout);
+            }
            const uint64 PresentOrdinal=++PresentCallCount;
+            if(bNonBlockingStartupPresent)
+            {
+                ++StartupNonBlockingPresentCalls;
+                if(PresentResult==DXGI_ERROR_WAS_STILL_DRAWING)++StartupNonBlockingPresentSkips;
+                else if(FAILED(PresentResult))
+                    std::fprintf(stderr,"[PresentStartupError] hr=0x%08lx\n",
+                        static_cast<unsigned long>(PresentResult));
+            }
            if(PresentOrdinal<=3)
-                std::printf("[PresentProfile] ordinal=%llu at_ms=%.3f sync=%u call_ms=%.3f gpu_tiles_submitted=%llu gpu_batches=%llu\n",
+                std::printf("[PresentProfile] ordinal=%llu at_ms=%.3f sync=%u flags=0x%08x mode=%s hr=0x%08lx call_ms=%.3f gpu_tiles_submitted=%llu gpu_batches=%llu\n",
                    static_cast<unsigned long long>(PresentOrdinal),
                    std::chrono::duration<double,std::milli>(PresentCallEnd-ColdProcessStart).count(),
-                    PresentSyncInterval,std::chrono::duration<double,std::milli>(PresentCallEnd-PresentCallStart).count(),
+                    PresentSyncInterval,PresentFlags,bNonBlockingStartupPresent?"startup-do-not-wait":"normal",
+                    static_cast<unsigned long>(PresentResult),
+                    std::chrono::duration<double,std::milli>(PresentCallEnd-PresentCallStart).count(),
                    static_cast<unsigned long long>(GpuFarTilesSubmitted),
                    static_cast<unsigned long long>(GpuFarSubmissionBatches));
+            if(bDeferGpuUntilCpuFrame&&!StreamingBatches.empty())
+            {
+                bCpuReadyFramePresentedBeforeGpu=true;
+                std::printf("[CpuReadyFrameBeforeGpu] at_ms=%.3f uploaded_cpu_batches=%zu next_action=submit_gpu_far\n",
+                    std::chrono::duration<double,std::milli>(PresentCallEnd-ColdProcessStart).count(),
+                    StreamingBatches.size());
+                std::fflush(stdout);
+            }
            UpdateTitle();
            const auto DrawEnd=std::chrono::steady_clock::now();
            const double CpuPumpMs=std::chrono::duration<double,std::milli>(CpuPumpEnd-DrawStart).count();
@@ -5230,13 +5427,22 @@ int main(int Argc,char** Argv)
        const std::string_view Arg=Argv[I];
        if(Arg.starts_with("--capture="))App.CapturePath=std::string(Arg.substr(10));
        else if(Arg=="--no-vsync")App.bDisableVSync=true;
+        else if(Arg=="--blocking-startup-present")App.bForceBlockingStartupPresent=true;
+        else if(Arg=="--cpu-frame-before-gpu")App.bPresentCpuFrameBeforeGpu=true;
+        else if(Arg=="--hold-cpu-frame-during-gpu-warmup")
+            App.bHoldCpuFrameDuringGpuWarmup=true,App.bPresentCpuFrameBeforeGpu=true;
+        else if(Arg=="--gpu-skirt-draw-audit")App.bAuditGpuSkirtDrawArgs=true;
+        else if(Arg=="--gpu-skirts-disabled")App.bGpuSkirtEmitEnabled=false;
+        else if(Arg=="--gpu-connectors-disabled")App.bGpuConnectorRepairEnabled=false;
        else if(Arg=="--profile-attribution")App.bProfileAttribution=true;
        else if(Arg=="--profile-detail-band")bProfileDetailBand=true;
        else if(Arg=="--compare-evaluators")App.bCompareEvaluators=true;
        else if(Arg=="--parity-only")App.bParityOnly=true;
        else if(Arg=="--density-hash")App.bDensityHash=true;
        else if(Arg=="--shape-world")App.bShapeWorld=true;
-        else if(Arg=="--gpu-far")App.bGpuFarEnabled=true,App.bShapeWorld=true,App.bForceGenerate=true;
+        else if(Arg=="--gpu-far")
+            App.bGpuFarEnabled=true,App.bShapeWorld=true,App.bForceGenerate=true,
+            App.bPresentCpuFrameBeforeGpu=true,App.bHoldCpuFrameDuringGpuWarmup=true;
        else if(Arg=="--cpu-far")App.bGpuFarEnabled=false,App.bShapeWorld=true,App.bForceGenerate=true;
        else if(Arg=="--gpu-far-audit")
        {App.bGpuFarEnabled=true;App.bGpuFarAuditEnabled=true;App.bShapeWorld=true;App.bForceGenerate=true;App.bQuitAfterClipmap=true;}
@@ -5316,6 +5522,20 @@ int main(int Argc,char** Argv)
        else if(Arg.starts_with("--capture-tile-limit="))
            App.CaptureTileLimit=std::clamp(std::atoi(
                std::string(Arg.substr(std::string_view("--capture-tile-limit=").size())).c_str()),1,4096);
+        else if(Arg.starts_with("--capture-only-lod="))
+            App.CaptureOnlyLod=std::clamp(std::atoi(
+                std::string(Arg.substr(std::string_view("--capture-only-lod=").size())).c_str()),0,8);
+        else if(Arg.starts_with("--capture-only-tile="))
+        {
+            const std::string Values(Arg.substr(std::string_view("--capture-only-tile=").size()));
+            if(std::sscanf(Values.c_str(),"%d,%d,%d,%d",&App.CaptureOnlyTileX,
+                &App.CaptureOnlyTileY,&App.CaptureOnlyTileZ,&App.CaptureOnlyTileLod)==4)
+                App.bCaptureOnlyTile=true;
+            else {std::fprintf(stderr,"Invalid --capture-only-tile value: %s\n",Values.c_str());return 2;}
+        }
+        else if(Arg.starts_with("--capture-output-slot="))
+            App.CaptureOutputSlot=std::max(-1,std::atoi(
+                std::string(Arg.substr(std::string_view("--capture-output-slot=").size())).c_str()));
        else if(Arg.starts_with("--benchmark-level="))
            BenchmarkLevel=std::atoi(std::string(Arg.substr(18)).c_str());
        else if(Arg.starts_with("--benchmark-count="))
@@ -5334,6 +5554,14 @@ int main(int Argc,char** Argv)
        }
        else if(!Arg.empty()&&Arg.front()!='-')App.DataPath=std::string(Arg);
    }
+    if(App.bCaptureOnlyTile&&App.bGpuFarEnabled)
+    {
+        // A single far-tile probe may have no CPU fallback batch to present.
+        // Let it submit directly instead of waiting for the normal CPU-first
+        // presentation gate that requires a nonempty CPU batch.
+        App.bPresentCpuFrameBeforeGpu=false;
+        App.bHoldCpuFrameDuringGpuWarmup=false;
+    }
    App.bProfileDetailBand=bProfileDetailBand;
    std::string Error;
    std::printf("[ColdTimeline] stage=owner_data_load begin_ms=%.3f\n",