diff --git a/Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h b/Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h index b5ca881..4eae1de 100644 --- a/Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h +++ b/Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h @@ -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 worldInput; // Sparse exception bitset: bit 1 forces route-protected air. Ordinary rock/air diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldViewer.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldViewer.ps1 index 5478e05..da5fe7f 100644 --- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldViewer.ps1 +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldViewer.ps1 @@ -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" diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep8Captures.py b/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep8Captures.py new file mode 100644 index 0000000..3ecf5cd --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep8Captures.py @@ -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}") diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep8LodIsolations.py b/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep8LodIsolations.py new file mode 100644 index 0000000..2b6aa0e --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/CompareStep8LodIsolations.py @@ -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}") diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp b/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp index 2f81644..01d7a27 100644 --- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp @@ -26,6 +26,7 @@ #include #include #include +#include #include #include @@ -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 vertices; + std::vector skirtVertices; std::vector cases; std::vector fieldDensity; ComPtr vertexBuffer, indexBuffer, indirectArgs; ComPtr connectorVertexBuffer,connectorIndexBuffer; ComPtr startTimestamp,endTimestamp,disjointTimestamp; std::array drawArgs{}; + std::array 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 kernelMs{}; + std::array kernelMs{}; std::uint64_t fieldDistanceEvaluations=0; std::uint32_t maxFieldDistanceEvaluationsPerSample=0; double wallMs = 0.0; @@ -506,6 +513,7 @@ namespace struct FarOutputChunk { ComPtr vertices, indices, args; + ComPtr surfaceVertexSrv; ComPtr vertexUav, indexUav, argsUav; ComPtr connectorVertices,connectorIndices; ComPtr connectorVertexUav,connectorIndexUav; @@ -528,6 +536,9 @@ namespace ComPtr connectorRingSrv; ComPtr densityUav,nearNormalDensityUav,caseUav,localOffsetUav,groupCountUav,groupOffsetUav; ComPtr connectorRingUav; + ComPtr skirtScratch; + ComPtr 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 immediateContext,deferredContext; ComPtr fieldEval, classify, fieldClassify, scan, emit, noiseProbe; ComPtr connectorRings,connectorPanels; + ComPtr classifyFarSkirts,scanFarSkirts,emitFarSkirts; ComPtr 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(std::chrono::steady_clock::now()-PipelineStart).count()); } @@ -783,17 +801,18 @@ namespace farOutputCapacitySlots-chunkIndex*farOutputSlotsPerChunk); const std::uint64_t vertexBytes=static_cast(slots)*farMaxVerticesPerTile*sizeof(GpuVertex); const std::uint64_t indexBytes=static_cast(slots)*farMaxVerticesPerTile*sizeof(std::uint32_t); - const std::uint64_t argsBytes=static_cast(slots)*10u*sizeof(std::uint32_t); + const std::uint64_t argsBytes=static_cast(slots)*11u*sizeof(std::uint32_t); if(vertexBytes>std::numeric_limits::max()||indexBytes>std::numeric_limits::max()) throw std::runtime_error("far output chunk exceeds the D3D11 buffer size limit"); FarOutputChunk chunk; chunk.vertices=RawMeshBuffer(static_cast(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(indexBytes), D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_INDEX_BUFFER); chunk.args=RawMeshBuffer(static_cast(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(vertexBytes/sizeof(std::uint32_t))); + chunk.surfaceVertexSrv=RawSrv(chunk.vertices.Get(),static_cast(vertexBytes/sizeof(std::uint32_t))); chunk.indexUav=RawUav(chunk.indices.Get(),static_cast(indexBytes/sizeof(std::uint32_t))); chunk.argsUav=RawUav(chunk.args.Get(),static_cast(argsBytes/sizeof(std::uint32_t))); const std::uint64_t connectorVertexBytes=static_cast(slots) @@ -885,6 +904,9 @@ namespace scratch.localOffsets=Structured(nullptr,cells,D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_SHADER_RESOURCE); scratch.groupCounts=Structured(nullptr,groups,D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_SHADER_RESOURCE); scratch.groupOffsets=Structured(nullptr,groups,D3D11_BIND_UNORDERED_ACCESS|D3D11_BIND_SHADER_RESOURCE); + scratch.skirtScratchCapacity=std::max(1,farMaxVerticesPerTile); + scratch.skirtScratch=Structured(nullptr,scratch.skirtScratchCapacity, + D3D11_BIND_UNORDERED_ACCESS); scratch.signSrv=Srv(scratch.signs.Get(),1); scratch.triSrv=Srv(scratch.tri.Get(),256u*16u); scratch.densitySrv=Srv(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 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 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(tile.dim) || farField->cells != static_cast(tile.cells) @@ -1153,36 +1189,49 @@ namespace } ComPtr args,vertices,indices; ComPtr argsUav,vertexUav,indexUav; + ComPtr surfaceVertexSrv; ComPtr connectorVertices,connectorIndices; ComPtr connectorVertexUav,connectorIndexUav; + ComPtr skirtScratchBuffer; + ComPtr skirtScratchUav; const std::uint32_t connectorPanelCapacity=pooledChunk?FarPanelVertexCapacity: - std::max(1,farField?farField->connectorVertexCount:0); + std::max(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(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(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(maxVertices*sizeof(GpuVertex)/sizeof(std::uint32_t))); indexUav=RawUav(indices.Get(),static_cast(maxVertices)); + surfaceVertexSrv=RawSrv(vertices.Get(),static_cast(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(nullptr, + static_cast(maxVertices),D3D11_BIND_UNORDERED_ACCESS); + skirtScratchUav=Uav(skirtScratchBuffer.Get(),static_cast(maxVertices)); + } } D3D11_BUFFER_DESC cbDesc{}; cbDesc.ByteWidth = sizeof(TileConstants); @@ -1213,7 +1262,9 @@ namespace const FarConnectorConstants connectorConstants{ static_cast(farField->connectorPanels.size()), farField->connectorRingCount,farField->connectorVertexCount,connectorPanelCapacity, - static_cast(farField->worldInput->shapes.size()),{0,0,0}}; + static_cast(farField->worldInput->shapes.size()), + static_cast(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 startQuery, endQuery, disjointQuery; ComPtr fieldStartQuery,fieldEndQuery; - std::array,6> kernelStartQueries,kernelEndQueries; + std::array,9> kernelStartQueries,kernelEndQueries; if(captureGpuTiming) { Check(device->CreateQuery(×tampDesc, &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(readbackCapacity * sizeof(std::uint32_t))); - auto argsStaging = Staging(10 * sizeof(std::uint32_t)); - auto connectorVertexStaging=farField&&farField->connectorVertexCount - ?Staging(farField->connectorVertexCount*sizeof(GpuVertex)):ComPtr{}; - auto connectorIndexStaging=farField&&farField->connectorVertexCount - ?Staging(farField->connectorVertexCount*sizeof(std::uint32_t)):ComPtr{}; + 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{}; + auto connectorIndexStaging=hasConnectorOutput + ?Staging(connectorPanelCapacity*sizeof(std::uint32_t)):ComPtr{}; 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;stageconnectorVertexCount==0)continue; + if((stage==5&&farField->connectorVertexCount==0) + ||(stage>=6&&farField->visibleSkirtFaceMask==0))continue; const UINT64 stageStart=ReadQuery(kernelStartQueries[stage].Get(), "GetData(kernel stage start)"); const UINT64 stageEnd=ReadQuery(kernelEndQueries[stage].Get(), @@ -1505,18 +1602,23 @@ namespace static_cast(mapped.pData) + cellCount); context->Unmap(caseStaging.Get(), 0); Check(context->Map(argsStaging.Get(), 0, D3D11_MAP_READ, 0, &mapped), "Map(args)"); - std::array allArgs{}; - std::copy_n(static_cast(mapped.pData),10,allArgs.begin()); + std::array allArgs{}; + std::copy_n(static_cast(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(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(result.drawArgs[0], static_cast(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(mapped.pData); + const std::uint32_t panelOutputCount=std::min( + 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(mapped.pData); - for(std::uint32_t i=0;iconnectorVertexCount;++i) + for(std::uint32_t i=0;iUnmap(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& desired) + { + const std::int64_t extent=static_cast(32)<(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(32)<(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& gpuSkirtVertices,std::uint8_t visibleFaceMask) + { + using FQuantizedVertex=std::array; + using FTriangleKey=std::array; + std::map cpuTriangles,gpuTriangles; + std::vector cpuVertices,gpuVertices; + std::vector> cpuGeometry,gpuGeometry; + const auto makeVertex=[](double x,double y,double z,double nx,double ny,double nz) + { + GpuVertex vertex{}; + vertex.position[0]=static_cast(x); + vertex.position[1]=static_cast(y); + vertex.position[2]=static_cast(z); + vertex.normal[0]=static_cast(nx); + vertex.normal[1]=static_cast(ny); + vertex.normal[2]=static_cast(nz); + return vertex; + }; + const auto triangleKey=[](const GpuVertex* vertices) + { + std::array points{}; + for(std::size_t i=0;i<3;++i) + points[i]={static_cast(std::llround(vertices[i].position[0])), + static_cast(std::llround(vertices[i].position[1])), + static_cast(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=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+2firstfirst)) + {result.quantizedKeyDifferenceTriangles+=cpuIt->second;++cpuIt;} + else if(cpuIt==cpuTriangles.end()||gpuIt->firstfirst) + {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(a.position[0])-b.position[0]; + const double dy=static_cast(a.position[1])-b.position[1]; + const double dz=static_cast(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 matched(gpuGeometry.size(),false); + for(const auto& cpuTriangle:cpuGeometry) + { + std::size_t best=gpuGeometry.size();double bestError=std::numeric_limits::infinity(); + for(std::size_t candidate=0;candidate( + 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(vertex.position[0])/25.0, + static_cast(vertex.position[1])/25.0, + static_cast(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::infinity();const GpuVertex* nearest=nullptr; + for(const GpuVertex& gpu:gpuVertices) + {const double distance=vertexDistance(cpu,gpu);if(distancemaxCpuToGpu){maxCpuToGpu=best;cpuWorst=&cpu;gpuNearest=nearest;} + } + for(const GpuVertex& gpu:gpuVertices) + { + double best=std::numeric_limits::infinity();const GpuVertex* nearest=nullptr; + for(const GpuVertex& cpu:cpuVertices) + {const double distance=vertexDistance(gpu,cpu);if(distancemaxGpuToCpu){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(&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(&hwFirst.skirtVertices[i]); + const std::uint32_t* b=reinterpret_cast(&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(&hwFirst.skirtVertices[i]); + const std::uint32_t* b=reinterpret_cast(&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(fieldSignDiff), static_cast(fieldCaseLatticeSignDiff), static_cast(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(farInput.visibleSkirtFaceMask), + static_cast(skirtAudit.cpuTriangles), + static_cast(skirtAudit.gpuTriangles), + static_cast(skirtAudit.topologyDifferenceTriangles), + static_cast(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(farInput.visibleSkirtFaceMask), + static_cast(warpSkirtAudit.cpuTriangles), + static_cast(warpSkirtAudit.gpuTriangles), + static_cast(warpSkirtAudit.topologyDifferenceTriangles), + static_cast(warpSkirtAudit.quantizedKeyDifferenceTriangles), + warpSkirtAudit.meanVertexDistanceCm,warpSkirtAudit.maxVertexDistanceCm, + warpSkirtAudit.vertexBoundCm,warp.connectorDrawArgs[5],warpSkirtAudit.pass?"PASS":"FAIL"); + } const std::size_t expectedGpuOutputCount=static_cast(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(noiseRepeat.differingBits), noiseWarp.mean, noiseWarp.maximum, static_cast(noiseWarp.differingBits)); - const std::vector selected = SelectFarTiles(world, config, camera,wholeFar,wholeNear); + std::vector selected = skirtAuditOnly + ?std::vector{}: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 desiredSet(desired.begin(),desired.end()); + selected.clear(); + for(const TileKey& key:desired) + { + if(key.lodconfig.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 gpuKernelStageSum{}; + std::array 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(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(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(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(selection.skirtFaceMask), + static_cast(exposedTriangles)); + } const auto inputStart = Clock::now(); TileData tile = MakeTileData(world, key, config, *cpuMesh, true,farInput); const double inputWallMs = std::chrono::duration(Clock::now() - inputStart).count(); @@ -2359,11 +2763,14 @@ namespace &&gpu1.vertices.size() == static_cast(tile.expectedVertexCount)+farInput->connectorVertexCount && gpu2.vertices.size() == gpu1.vertices.size() && warpResult.vertices.size() == static_cast(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(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(testedTiles),gpuKernelStageSum[0],gpuKernelStageSum[1], gpuKernelStageSum[2],gpuKernelStageSum[3],gpuKernelStageSum[4],gpuKernelStageSum[5], + gpuKernelStageSum[6],gpuKernelStageSum[7],gpuKernelStageSum[8], static_cast(gpuFieldDistanceEvaluations), static_cast(gpuFieldSamples), gpuFieldSamples?static_cast(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;ivertices[i].normal[c]=result.vertices[i].normal[c]; } } - auditReadback->indices.resize(result.vertices.size()); + for(std::size_t i=0;ivertices[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) diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8CaptureMatrix.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8CaptureMatrix.ps1 new file mode 100644 index 0000000..3dab6d4 --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8CaptureMatrix.ps1 @@ -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 + } + } +} diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8ColdInterleaved.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8ColdInterleaved.ps1 new file mode 100644 index 0000000..be23787 --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8ColdInterleaved.ps1 @@ -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 + } +} diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8PresentControls.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8PresentControls.ps1 new file mode 100644 index 0000000..e948bb3 --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8PresentControls.ps1 @@ -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 + } +} diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8PresentVisibilityControls.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8PresentVisibilityControls.ps1 new file mode 100644 index 0000000..233e825 --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunStep8PresentVisibilityControls.ps1 @@ -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 + } +} diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/BuildFarConnectorRingsCS.cso b/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/BuildFarConnectorRingsCS.cso index 27d2ca6..7a26919 100644 Binary files a/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/BuildFarConnectorRingsCS.cso and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/BuildFarConnectorRingsCS.cso differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/EmitFarConnectorPanelsCS.cso b/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/EmitFarConnectorPanelsCS.cso index 22db721..dbaa94e 100644 Binary files a/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/EmitFarConnectorPanelsCS.cso and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/EmitFarConnectorPanelsCS.cso differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/ScanGroupsCS.cso b/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/ScanGroupsCS.cso index a57a694..c2d801d 100644 Binary files a/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/ScanGroupsCS.cso and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/precompiled/ScanGroupsCS.cso differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_capture_comparison.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_capture_comparison.md new file mode 100644 index 0000000..3389730 --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_capture_comparison.md @@ -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% | diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_clamped_capture_comparison.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_clamped_capture_comparison.md new file mode 100644 index 0000000..f44d340 --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_clamped_capture_comparison.md @@ -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% | diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_chamber.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_chamber.png new file mode 100644 index 0000000..7674db3 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_chamber.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_far.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_far.png new file mode 100644 index 0000000..db19e92 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_far.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_passage.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_passage.png new file mode 100644 index 0000000..84078c1 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_passage.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_passage_floor_closeup.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_passage_floor_closeup.png new file mode 100644 index 0000000..0a0eac3 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_passage_floor_closeup.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw0.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw0.png new file mode 100644 index 0000000..ea00fb6 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw0.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw180.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw180.png new file mode 100644 index 0000000..84d0a2a Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw180.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw270.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw270.png new file mode 100644 index 0000000..61ee3b1 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw270.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw90.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw90.png new file mode 100644 index 0000000..b793ef6 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_remote_yaw90.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_tunnel.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_tunnel.png new file mode 100644 index 0000000..aeacfa6 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_tunnel.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_tunnel_bend.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_tunnel_bend.png new file mode 100644 index 0000000..96e2172 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_cpu_tunnel_bend.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_chamber.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_chamber.png new file mode 100644 index 0000000..dc89240 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_chamber.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_far.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_far.png new file mode 100644 index 0000000..1d39010 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_far.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_passage.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_passage.png new file mode 100644 index 0000000..4345c1d Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_passage.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_passage_floor_closeup.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_passage_floor_closeup.png new file mode 100644 index 0000000..6620445 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_passage_floor_closeup.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_protected_remote_yaw0.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_protected_remote_yaw0.png new file mode 100644 index 0000000..d19264e Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_protected_remote_yaw0.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_protected_remote_yaw90.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_protected_remote_yaw90.png new file mode 100644 index 0000000..a7e55a0 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_protected_remote_yaw90.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw0.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw0.png new file mode 100644 index 0000000..9b64141 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw0.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw180.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw180.png new file mode 100644 index 0000000..e6b71c7 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw180.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw270.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw270.png new file mode 100644 index 0000000..d87384f Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw270.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90.png new file mode 100644 index 0000000..a7e55a0 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90_direct_audit.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90_direct_audit.png new file mode 100644 index 0000000..a7e55a0 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90_direct_audit.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90_drawargs.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90_drawargs.png new file mode 100644 index 0000000..a7e55a0 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90_drawargs.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90_drawargs_full.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90_drawargs_full.png new file mode 100644 index 0000000..a7e55a0 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_remote_yaw90_drawargs_full.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_tunnel.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_tunnel.png new file mode 100644 index 0000000..d2e6729 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_tunnel.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_tunnel_bend.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_tunnel_bend.png new file mode 100644 index 0000000..3866f6a Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_gpu_tunnel_bend.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_lod_isolation_comparison.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_lod_isolation_comparison.md new file mode 100644 index 0000000..fb5314b --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step8_lod_isolation_comparison.md @@ -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 | diff --git a/Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl b/Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl index 0286e64..9185fa9 100644 --- a/Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl +++ b/Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl @@ -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 FarConnectorRingOutput : register(u7); RWByteAddressBuffer FarConnectorDrawArgs : register(u4); RWByteAddressBuffer FarConnectorVertexBytes : register(u5); RWByteAddressBuffer FarConnectorIndexBytes : register(u6); +RWStructuredBuffer FarSkirtScratch : register(u0); +ByteAddressBuffer FarSurfaceVertexInput : register(t19); RWStructuredBuffer 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=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 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;edgeVertexIdFarConnectorPanelVertexCapacity) + { + 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; }; diff --git a/Standalone/Source/StandaloneViewer.cpp b/Standalone/Source/StandaloneViewer.cpp index 320e1a9..4f49cc5 100644 --- a/Standalone/Source/StandaloneViewer.cpp +++ b/Standalone/Source/StandaloneViewer.cpp @@ -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 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(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.LevelFullResClipLevels ?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(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(Eligible.size()); + if(bCaptureOnlyTile&&CaptureOutputSlot>=0) + { + GpuFarOutputCapacitySlots=std::max(GpuFarOutputCapacitySlots, + static_cast(CaptureOutputSlot+1)); + GpuFarNextOutputSlot=static_cast(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.LevelFullResClipLevels;})); 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(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( 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 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(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(StartupNonBlockingPresentCalls), + static_cast(StartupNonBlockingPresentSkips), + static_cast(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 "<(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(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(ArgsMapped.pData); + const uint32 ExpectedStartIndex=(Tile.IndirectArgsByteOffset + /(11u*sizeof(uint32)))*32768u; + ++GpuSkirtDrawAuditTiles; + GpuSkirtDrawAuditSumIndices+=Args[5]; + GpuSkirtDrawAuditMaxIndices=std::max(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(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(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(GpuSkirtDrawAuditTiles), + static_cast(GpuSkirtDrawAuditZeroCounts), + static_cast(GpuSkirtDrawAuditOverflow), + static_cast(GpuSkirtDrawAuditBadStartIndex), + static_cast(GpuSkirtDrawAuditSumIndices), + static_cast(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( + 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(PresentCallEnd-ColdProcessStart).count(), + std::chrono::duration(PresentCallEnd-PresentCallStart).count(), + PresentSyncInterval,static_cast(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(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(PresentOrdinal), std::chrono::duration(PresentCallEnd-ColdProcessStart).count(), - PresentSyncInterval,std::chrono::duration(PresentCallEnd-PresentCallStart).count(), + PresentSyncInterval,PresentFlags,bNonBlockingStartupPresent?"startup-do-not-wait":"normal", + static_cast(PresentResult), + std::chrono::duration(PresentCallEnd-PresentCallStart).count(), static_cast(GpuFarTilesSubmitted), static_cast(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(PresentCallEnd-ColdProcessStart).count(), + StreamingBatches.size()); + std::fflush(stdout); + } UpdateTitle(); const auto DrawEnd=std::chrono::steady_clock::now(); const double CpuPumpMs=std::chrono::duration(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",