diff --git a/Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h b/Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h index 42a8019..4582906 100644 --- a/Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h +++ b/Standalone/Core/include/VoxelForgeCore/VoxelForgeCore.h @@ -215,12 +215,18 @@ namespace VoxelForgeCore // Shape IDs preserve the CPU spatial-bin iteration order. std::vector sampleShapeRanges; std::vector sampleShapeIds; - // Bit 1 means air, matching the CPU density sign rule (density >= 0). - std::vector rockAirSigns; - // Critical route air is a subset of rockAirSigns and is always authoritative. + // Sparse exception bitset: bit 1 forces route-protected air. Ordinary rock/air + // classification is computed from the GPU field and is never uploaded by the CPU. std::vector protectedAir; + // Sparse exception bitset for exact solid overlays/thin route-adjacent features. + std::vector protectedSolid; + std::vector protectedThin; std::vector stratumBoundaries; std::uint64_t protectedSampleCount = 0; + std::uint64_t protectedSolidSampleCount = 0; + std::uint64_t protectedSealSampleCount = 0; + std::uint64_t protectedThinSampleCount = 0; + std::uint64_t protectedSolidOverlaySampleCount = 0; }; struct EditBrush @@ -312,6 +318,7 @@ namespace VoxelForgeCore std::vector farConnectorCandidates(Vec3i minVoxel, Vec3i maxVoxel) const; std::optional exportFarFieldTile(const TileKey& key, std::string& error) const; + bool isFarFieldTileEligible(const TileKey& key) const; bool auditDetailPassing() const; bool auditGridStairIntegrity() const; diff --git a/Standalone/Core/src/ShapeWorldInternal.h b/Standalone/Core/src/ShapeWorldInternal.h index 19d2e43..c9f8eaf 100644 --- a/Standalone/Core/src/ShapeWorldInternal.h +++ b/Standalone/Core/src/ShapeWorldInternal.h @@ -501,8 +501,9 @@ public: {return std::abs(A.z-B.z)<0.01f;}),Out.stratumBoundaries.end()); } Out.sampleShapeRanges.resize(SampleCount); - Out.rockAirSigns.assign((SampleCount+31u)/32u,0u); Out.protectedAir.assign((SampleCount+31u)/32u,0u); + Out.protectedSolid.assign((SampleCount+31u)/32u,0u); + Out.protectedThin.assign((SampleCount+31u)/32u,0u); size_t SampleIndex=0; for(int32 Z=-1;Z<=Cells+1;++Z) @@ -512,13 +513,53 @@ public: const int32 WX=Origin.X+X*Step,WY=Origin.Y+Y*Step,WZ=Origin.Z+Z*Step; if(WX>=Min[0]&&WX<=Max[0]&&WY>=Min[1]&&WY<=Max[1]&&WZ>=Min[2]&&WZ<=Max[2]) {Error="tile lattice intersects the near analytic/detail region; this export is far-field only";return false;} - const float Density=DensityAtLattice(WX,WY,WZ); - if(Density>=0.0f)Out.rockAirSigns[SampleIndex>>5]|=1u<<(SampleIndex&31u); if(IsInsideCriticalAir(WX,WY,WZ)) { Out.protectedAir[SampleIndex>>5]|=1u<<(SampleIndex&31u); ++Out.protectedSampleCount; } + else if(IsNearCriticalAirBoundary(WX,WY,WZ,4.0)) + { + // Preserve exact signs within roughly one far-field lattice step of + // route air. This catches thin walls/floors at the route boundary + // without evaluating the ordinary far field on the CPU. + const float CriticalDensity=AnalyticDensityAt(static_cast(WX), + static_cast(WY),static_cast(WZ)); + Out.protectedThin[SampleIndex>>5]|=1u<<(SampleIndex&31u); + if(CriticalDensity>=0.0f)Out.protectedAir[SampleIndex>>5]|=1u<<(SampleIndex&31u); + else + { + Out.protectedSolid[SampleIndex>>5]|=1u<<(SampleIndex&31u); + ++Out.protectedSolidSampleCount; + } + ++Out.protectedThinSampleCount; + } + else + { + bool bSealSample=false; + for(const VoxelForgeCore::FarFieldStratumBoundary& Boundary:Out.stratumBoundaries) + { + if(std::abs(static_cast(WZ)-Boundary.z)>0.5)continue; + // Only the exact lattice samples that lie on a seal height need a + // CPU sign override. The surrounding far field remains GPU-decided. + const float SealDensity=AnalyticDensityAt(static_cast(WX), + static_cast(WY),static_cast(WZ)); + if(SealDensity>=0.0f) + Out.protectedAir[SampleIndex>>5]|=1u<<(SampleIndex&31u); + else + Out.protectedSolid[SampleIndex>>5]|=1u<<(SampleIndex&31u); + ++Out.protectedSealSampleCount; + if(SealDensity<0.0f)++Out.protectedSolidSampleCount; + bSealSample=true; + break; + } + if(!bSealSample&&IsInsideSolidOverlay(WX,WY,WZ)) + { + Out.protectedSolid[SampleIndex>>5]|=1u<<(SampleIndex&31u); + ++Out.protectedSolidSampleCount; + ++Out.protectedSolidOverlaySampleCount; + } + } const auto Found=SpatialBins.find(PointBin(WX,WY,WZ)); if(Found==SpatialBins.end())continue; auto& Ids=PerSample[SampleIndex]; @@ -586,6 +627,19 @@ public: return true; } + bool IsFarFieldTileOutsideNearRegion(const FIntVector& Origin,int32 Cells,int32 Step) const + { + if(!FarBuilt||Cells<=0||Cells>32||Step<=0)return false; + const int32 Origins[3]={Origin.X,Origin.Y,Origin.Z}; + for(int32 Axis=0;Axis<3;++Axis) + { + const double First=std::ceil((static_cast(Min[Axis])-Origins[Axis])/Step); + const double Last=std::floor((static_cast(Max[Axis])-Origins[Axis])/Step); + if(First>Cells+1||Last< -1)return true; + } + return false; + } + bool DebugCriticalAirAt(double X,double Y,double Z) const { return IsInsideCriticalAir(X,Y,Z); } bool DebugSolidOverlayAt(double X,double Y,double Z) const { return IsInsideSolidOverlay(X,Y,Z); } @@ -3336,6 +3390,27 @@ private: return false; } + bool IsNearCriticalAirBoundary(int32 X,int32 Y,int32 Z,double Margin) const + { + const auto It=CriticalAirBins.find(PointBin(X,Y,Z)); + if(It==CriticalAirBins.end())return false; + for(uint32 Id:It->second) + { + const FPrimitive& P=Primitives[Id];if(!P.bCarve||!P.bCritical)continue; + if(XP.Max[0]+Margin + ||YP.Max[1]+Margin + ||ZP.Max[2]+Margin)continue; + float Distance=ShapeDistance(P,X,Y,Z); + if(P.bHasFloorClip) + { + const float FloorDistance=static_cast(FloorAt(P,X,Y)-Z); + Distance=SmoothMaxDistance(Distance,FloorDistance,static_cast(P.FloorBlendRadius)); + } + if(std::abs(Distance)<=Margin)return true; + } + return false; + } + bool IsInsideSolidOverlay(double X,double Y,double Z) const { const auto It=SpatialBins.find(PointBin(X,Y,Z)); diff --git a/Standalone/Core/src/VoxelForgeCore.cpp b/Standalone/Core/src/VoxelForgeCore.cpp index 5f4f401..46e7b4f 100644 --- a/Standalone/Core/src/VoxelForgeCore.cpp +++ b/Standalone/Core/src/VoxelForgeCore.cpp @@ -751,6 +751,17 @@ namespace VoxelForgeCore return result; } + bool World::isFarFieldTileEligible(const TileKey& key) const + { + if(!impl_||!impl_->owner||!impl_->shape||!impl_->shape->IsBuilt() + ||!impl_->shape->HasFarWorld()||key.lodconfiguration.fullResolutionLodCount + ||key.lod>8)return false; + const int32 extent=CHUNK_SIZE<owner->CoarseTileCells); + const int32 step=std::max(1,extent/cells); + return impl_->shape->IsFarFieldTileOutsideNearRegion(ToNative(key).OriginVoxels(),cells,step); + } + bool World::auditDetailPassing() const { return impl_&&impl_->shape&&impl_->shape->IsDetailAuditPassing(); } bool World::auditGridStairIntegrity() const { return impl_&&impl_->shape&&impl_->shape->AuditGridStairIntegrity(); } diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldPort.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldPort.ps1 index 015edba..82eb989 100644 --- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldPort.ps1 +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldPort.ps1 @@ -1,15 +1,15 @@ [CmdletBinding()] -param([switch]$SkipCore) +param([switch]$SkipCore,[string]$OutputName='VoxelForgeGpuFarAgreement.exe') $ErrorActionPreference = 'Stop' $HostProcess = Get-Process -Id $PID $HostProcess.PriorityClass = 'BelowNormal' $HostProcess.ProcessorAffinity = [IntPtr]0xFF00 $Standalone = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..')) -$BuildDir = Join-Path $Standalone 'Build\VoxelForgeGpuFieldPort' +$BuildDir = Join-Path $Standalone 'Build\VoxelForgeGpuFarAgreement' $WrapperDir = Join-Path $BuildDir 'wrappers' $BinDir = Join-Path $Standalone 'bin' -$Output = Join-Path $BinDir 'VoxelForgeGpuFieldPort.exe' +$Output = Join-Path $BinDir $OutputName $CoreLib = Join-Path $BuildDir 'VoxelForgeGpuFieldPortCore.lib' $VcVars = 'E:\Program Files\Microsoft Visual Studio\2022\Community\VC\Auxiliary\Build\vcvars64.bat' diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldViewer.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldViewer.ps1 index cd1207c..e2f66ee 100644 --- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldViewer.ps1 +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/BuildGpuFieldViewer.ps1 @@ -1,6 +1,6 @@ [CmdletBinding()] param( - [string]$OutputName = 'VoxelForgeGpuFieldViewer.exe', + [string]$OutputName = 'VoxelForgeGpuFarHybridViewer.exe', [ValidateSet('All','Core')][string]$Target = 'All' ) @@ -9,7 +9,7 @@ $BuildHostProcess = Get-Process -Id $PID $BuildHostProcess.PriorityClass = 'BelowNormal' $BuildHostProcess.ProcessorAffinity = [IntPtr]0xFF00 $Root = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..')) -$BuildDir = Join-Path $Root 'Build\VoxelForgeGpuFieldViewer' +$BuildDir = Join-Path $Root 'Build\VoxelForgeGpuFarHybridViewer' $BinDir = Join-Path $Root 'bin' [IO.Directory]::CreateDirectory($BuildDir) | Out-Null [IO.Directory]::CreateDirectory($BinDir) | Out-Null @@ -87,7 +87,8 @@ $IncludeArgs = @( ('/I"{0}"' -f (Join-Path $Root 'Compat')), ('/I"{0}"' -f (Join-Path $Root 'Source')), ('/I"{0}"' -f (Join-Path $Root 'Source\VoxelForge\Public')), - ('/I"{0}"' -f (Join-Path $Root 'Source\VoxelForge\Private')) + ('/I"{0}"' -f (Join-Path $Root 'Source\VoxelForge\Private')), + ('/I"{0}"' -f $PSScriptRoot) ) $CommonFlags = @('/nologo','/std:c++20','/O2','/arch:AVX2','/fp:precise','/EHsc','/MP8','/D_CRT_SECURE_NO_WARNINGS') + $IncludeArgs $CoreRsp = Join-Path $BuildDir 'core-compile.rsp' @@ -112,10 +113,16 @@ if ($Target -eq 'All') { } $ViewerRsp = Join-Path $BuildDir 'viewer-link.rsp' +$GpuViewerObject = Join-Path $BuildDir 'GpuFieldPortViewer.obj' +$GpuViewerCompileRsp = Join-Path $BuildDir 'gpu-field-viewer-compile.rsp' +$GpuViewerCompileLines = $CommonFlags + @('/DVF_GPU_FIELD_PORT_NO_MAIN',('/Fo"{0}"' -f $GpuViewerObject),'/c', + ('"{0}"' -f (Join-Path $PSScriptRoot 'GpuFieldPort.cpp'))) +[IO.File]::WriteAllLines($GpuViewerCompileRsp,$GpuViewerCompileLines,[Text.UTF8Encoding]::new($false)) $ViewerRspLines = $CommonFlags + @('/Fe"{0}"' -f $OutputPath) + $LegacyObjects + @( + ('"{0}"' -f $GpuViewerObject), ('"{0}"' -f (Join-Path $Root 'Source\StandaloneViewer.cpp')), ('"{0}"' -f $CoreLibrary), - '/link','d3d11.lib','dxgi.lib','d3dcompiler.lib','user32.lib','gdi32.lib' + '/link','d3d11.lib','dxgi.lib','d3dcompiler.lib','user32.lib','gdi32.lib','psapi.lib' ) [IO.File]::WriteAllLines($ViewerRsp,$ViewerRspLines,[Text.UTF8Encoding]::new($false)) @@ -131,7 +138,7 @@ try { if ($Target -eq 'All') { # The standalone viewer retains a global compatibility build only for legacy comparison # and audits. Those objects are linked into the viewer executable, never into Core.lib. - $CommandLine = 'call "{0}" >nul && cl.exe @"{1}" && cl.exe @"{2}"' -f $VcVars,$LegacyRsp,$ViewerRsp + $CommandLine = 'call "{0}" >nul && cl.exe @"{1}" && cl.exe @"{2}" && cl.exe @"{3}"' -f $VcVars,$LegacyRsp,$GpuViewerCompileRsp,$ViewerRsp Invoke-FieldViewerBuildCommand $CommandLine 'standalone viewer compile/link step' } if ($Target -eq 'All' -and -not (Test-Path -LiteralPath $OutputPath -PathType Leaf)) { throw "Viewer was not produced: $OutputPath" } diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h b/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h new file mode 100644 index 0000000..7eaa417 --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h @@ -0,0 +1,28 @@ +#pragma once + +#include +#include +#include +#include +#include "VoxelForgeCore/VoxelForgeCore.h" + +struct ID3D11Device; +struct ID3D11DeviceContext; +struct ID3D11Buffer; +struct ID3D11Query; + +namespace VoxelForgeGpuFarViewer +{ + struct CpuVertex { float position[3]{}, normal[3]{}; }; + struct CpuMesh { std::vector vertices; std::vector indices; }; + // Returned D3D11 buffers carry one reference each; the caller releases them. + // Vertex records keep the 32-byte compute layout (position/normal/uv). The + // first 24 bytes match the standalone viewer's input layout. + bool BuildFarTile(ID3D11Device* device, ID3D11DeviceContext* context, + VoxelForgeCore::World& world, const VoxelForgeCore::TileKey& key, + const std::filesystem::path& shaderPath, + ID3D11Buffer** vertices, ID3D11Buffer** indices, ID3D11Buffer** indirectArgs, + ID3D11Query** completionEvent, + CpuMesh* auditReadback, + double* enqueueWallMs, double* gpuDispatchMs, std::string& error); +} diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp b/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp index ab1bd84..a601fce 100644 --- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/GpuFieldPort.cpp @@ -144,18 +144,21 @@ namespace std::vector signs; std::vector expectedCases; std::vector criticalCellMask; + std::vector thinFeatureCellMask; std::vector midpointIds; std::vector midpointRecords; std::vector cpuSurfaceVertices; std::uint32_t expectedVertexCount = 0; std::uint64_t criticalSamples = 0; std::uint64_t criticalCells = 0; + std::uint64_t thinFeatureCells = 0; }; struct GpuResult { std::vector vertices; std::vector cases; std::vector fieldDensity; + ComPtr vertexBuffer, indexBuffer, indirectArgs; std::array drawArgs{}; bool sequentialIndices = false; bool indexedIndirectArgsValid = false; @@ -287,7 +290,8 @@ namespace } TileData MakeTileData(World& world, const TileKey& key, const WorldConfiguration& config, - const TileMesh& cpuMesh, bool auditCritical) + const TileMesh& cpuMesh, bool auditCritical, + const FarFieldTileInput* farInput) { TileData t; t.key = key; @@ -309,7 +313,12 @@ namespace const float d = world.densityAtLattice(p); t.density[i] = d; if (d >= 0.0f) t.signs[i >> 5] |= 1u << (i & 31); - if (auditCritical && world.debugCriticalAirAt({double(p.x), double(p.y), double(p.z)})) + const bool routeException=farInput&&i/32protectedAir.size() + &&((farInput->protectedAir[i>>5]>>(i&31u))&1u)!=0; + const bool solidException=farInput&&i/32protectedSolid.size() + &&((farInput->protectedSolid[i>>5]>>(i&31u))&1u)!=0; + if (auditCritical && (routeException||solidException + ||world.debugCriticalAirAt({double(p.x), double(p.y), double(p.z)}))) ++t.criticalSamples; } @@ -325,6 +334,7 @@ namespace const std::uint32_t cellCount = t.cells * t.cells * t.cells; t.expectedCases.resize(cellCount); t.criticalCellMask.assign(cellCount, 0); + t.thinFeatureCellMask.assign(cellCount, 0); t.midpointIds.assign(cellCount, InvalidIndex); for (std::uint32_t z = 0; z < t.cells; ++z) for (std::uint32_t y = 0; y < t.cells; ++y) @@ -338,19 +348,35 @@ namespace t.expectedVertexCount += vertexCount; if (auditCritical) { - bool hasCritical = false; - for (int corner = 0; corner < 8 && !hasCritical; ++corner) + bool hasCritical = false,hasThinFeature=false; + for (int corner = 0; corner < 8; ++corner) { const Vec3d p{double(t.origin[0] + (static_cast(x) + Corner[corner][0]) * t.step), double(t.origin[1] + (static_cast(y) + Corner[corner][1]) * t.step), double(t.origin[2] + (static_cast(z) + Corner[corner][2]) * t.step)}; - hasCritical = world.debugCriticalAirAt(p); + const std::uint32_t sample=SampleIndex(t, + static_cast(x)+Corner[corner][0], + static_cast(y)+Corner[corner][1], + static_cast(z)+Corner[corner][2]); + const bool routeException=farInput&&sample/32protectedAir.size() + &&((farInput->protectedAir[sample>>5]>>(sample&31u))&1u)!=0; + const bool solidException=farInput&&sample/32protectedSolid.size() + &&((farInput->protectedSolid[sample>>5]>>(sample&31u))&1u)!=0; + const bool thinException=farInput&&sample/32protectedThin.size() + &&((farInput->protectedThin[sample>>5]>>(sample&31u))&1u)!=0; + hasCritical=hasCritical||routeException||solidException||world.debugCriticalAirAt(p); + hasThinFeature=hasThinFeature||thinException; } if (hasCritical) { t.criticalCellMask[cell] = 1; ++t.criticalCells; } + if (hasThinFeature) + { + t.thinFeatureCellMask[cell] = 1; + ++t.thinFeatureCells; + } } if (MCEdgeTable[mcCase] == 0) continue; @@ -431,11 +457,37 @@ namespace { ComPtr device; ComPtr context; - ComPtr fieldEval, classify, scan, emit, noiseProbe; + ComPtr fieldEval, classify, fieldClassify, scan, emit, noiseProbe; ComPtr memoryAdapter; std::string name; UINT64 timestampFrequency = 0; + void CompileShaders(const std::filesystem::path& shaderPath) + { + ComPtr code = CompileShader(shaderPath, "EvaluateFarFieldCS"); + Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &fieldEval), "CreateComputeShader(EvaluateFarFieldCS)"); + code = CompileShader(shaderPath, "ClassifyCS"); + Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &classify), "CreateComputeShader(ClassifyCS)"); + code = CompileShader(shaderPath, "ClassifyFarCS"); + Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &fieldClassify), "CreateComputeShader(ClassifyFarCS)"); + code = CompileShader(shaderPath, "ScanGroupsCS"); + Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &scan), "CreateComputeShader(ScanGroupsCS)"); + code = CompileShader(shaderPath, "EmitCS"); + Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &emit), "CreateComputeShader(EmitCS)"); + code = CompileShader(shaderPath, "NoiseProbeCS"); + Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &noiseProbe), "CreateComputeShader(NoiseProbeCS)"); + } + + void InitializeExternal(ID3D11Device* inDevice, ID3D11DeviceContext* inContext, + const std::filesystem::path& shaderPath) + { + device = inDevice; + context = inContext; + if (!device || !context) throw std::runtime_error("GPU far viewer requires a live D3D11 device/context"); + name = "viewer-device"; + CompileShaders(shaderPath); + } + void Initialize(const std::filesystem::path& shaderPath, bool warp) { D3D_FEATURE_LEVEL requested[] = {D3D_FEATURE_LEVEL_11_1, D3D_FEATURE_LEVEL_11_0}; @@ -462,16 +514,7 @@ namespace if (!warp) name = narrow; adapter.As(&memoryAdapter); } - ComPtr code = CompileShader(shaderPath, "EvaluateFarFieldCS"); - Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &fieldEval), "CreateComputeShader(EvaluateFarFieldCS)"); - code = CompileShader(shaderPath, "ClassifyCS"); - Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &classify), "CreateComputeShader(ClassifyCS)"); - code = CompileShader(shaderPath, "ScanGroupsCS"); - Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &scan), "CreateComputeShader(ScanGroupsCS)"); - code = CompileShader(shaderPath, "EmitCS"); - Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &emit), "CreateComputeShader(EmitCS)"); - code = CompileShader(shaderPath, "NoiseProbeCS"); - Check(device->CreateComputeShader(code->GetBufferPointer(), code->GetBufferSize(), nullptr, &noiseProbe), "CreateComputeShader(NoiseProbeCS)"); + CompileShaders(shaderPath); } bool QueryLocalMemory(std::uint64_t& usage, std::uint64_t& budget) const @@ -598,7 +641,8 @@ namespace } GpuResult Run(const TileData& tile, const FarFieldTileInput* farField = nullptr, - bool readback = true) + bool readback = true, bool retainOutputBuffers = false, + bool waitForGpu = true, bool useIndirectOutputCount = false) { using namespace std::chrono; const auto wallStart = Clock::now(); @@ -637,8 +681,8 @@ namespace result.d3dBufferBytes += std::max(1, farField->shapes.size()) * sizeof(FarFieldPrimitiveGpu) + farField->sampleShapeRanges.size() * sizeof(FarFieldSampleRange) + std::max(1, farField->sampleShapeIds.size()) * sizeof(std::uint32_t) - + farField->rockAirSigns.size() * sizeof(std::uint32_t) + farField->protectedAir.size() * sizeof(std::uint32_t) + + farField->protectedSolid.size() * sizeof(std::uint32_t) + tile.density.size() * sizeof(float) + sizeof(FarFieldConstants); if (readback) @@ -681,12 +725,12 @@ namespace auto farIds = farField ? Structured(farField->sampleShapeIds.empty()?&zeroId:farField->sampleShapeIds.data(), std::max(1,farField->sampleShapeIds.size()),D3D11_BIND_SHADER_RESOURCE) : ComPtr{}; - auto farSigns = farField ? Structured(farField->rockAirSigns.empty()?&zeroWord:farField->rockAirSigns.data(), - std::max(1,farField->rockAirSigns.size()),D3D11_BIND_SHADER_RESOURCE) - : ComPtr{}; auto farProtected = farField ? Structured(farField->protectedAir.empty()?&zeroWord:farField->protectedAir.data(), std::max(1,farField->protectedAir.size()),D3D11_BIND_SHADER_RESOURCE) : ComPtr{}; + auto farProtectedSolid = farField ? Structured(farField->protectedSolid.empty()?&zeroWord:farField->protectedSolid.data(), + std::max(1,farField->protectedSolid.size()),D3D11_BIND_SHADER_RESOURCE) + : ComPtr{}; auto signSrv = Srv(signs.Get(), tile.signs.size()); auto triSrv = Srv(triBuffer.Get(), triCount); auto densitySrv = Srv(density.Get(), tile.density.size()); @@ -696,10 +740,10 @@ namespace : ComPtr{}; auto farIdSrv = farField ? Srv(farIds.Get(),std::max(1,farField->sampleShapeIds.size())) : ComPtr{}; - auto farSignSrv = farField ? Srv(farSigns.Get(),std::max(1,farField->rockAirSigns.size())) - : ComPtr{}; auto farProtectedSrv = farField ? Srv(farProtected.Get(),std::max(1,farField->protectedAir.size())) : ComPtr{}; + auto farProtectedSolidSrv = farField ? Srv(farProtectedSolid.Get(),std::max(1,farField->protectedSolid.size())) + : ComPtr{}; auto midpointIdSrv = Srv(midpointIds.Get(), tile.midpointIds.size()); auto midpointRecordSrv = Srv(midpointRecords.Get(), std::max(1, tile.midpointRecords.size())); auto caseSrv = Srv(cases.Get(), cellCount); @@ -758,8 +802,8 @@ namespace context->CSSetShader(fieldEval.Get(),nullptr,0); ID3D11ShaderResourceView* fieldSrvs[15]{}; fieldSrvs[10]=farShapeSrv.Get();fieldSrvs[11]=farRangeSrv.Get(); - fieldSrvs[12]=farIdSrv.Get();fieldSrvs[13]=farSignSrv.Get(); - fieldSrvs[14]=farProtectedSrv.Get(); + fieldSrvs[12]=farIdSrv.Get();fieldSrvs[13]=farProtectedSrv.Get(); + fieldSrvs[14]=farProtectedSolidSrv.Get(); context->CSSetShaderResources(0,15,fieldSrvs); ID3D11UnorderedAccessView* fieldUavs[8]{};fieldUavs[7]=densityUav.Get(); context->CSSetUnorderedAccessViews(0,8,fieldUavs,nullptr); @@ -770,9 +814,11 @@ namespace context->CSSetUnorderedAccessViews(0,8,clearUavs,nullptr); context->End(fieldEndQuery.Get()); } - context->CSSetShader(classify.Get(), nullptr, 0); - ID3D11ShaderResourceView* classifySrvs[] = {signSrv.Get(), triSrv.Get()}; - context->CSSetShaderResources(0, 2, classifySrvs); + context->CSSetShader(farField ? fieldClassify.Get() : classify.Get(), nullptr, 0); + ID3D11ShaderResourceView* classifySrvs[16]{}; + classifySrvs[0]=signSrv.Get();classifySrvs[1]=triSrv.Get(); + if(farField)classifySrvs[15]=densitySrv.Get(); + context->CSSetShaderResources(0, 16, classifySrvs); ID3D11UnorderedAccessView* classifyUavs[] = {caseUav.Get(), localOffsetUav.Get(), groupCountUav.Get()}; context->CSSetUnorderedAccessViews(0, 3, classifyUavs, nullptr); context->Dispatch(groupCount, 1, 1); @@ -809,9 +855,11 @@ namespace auto caseStaging = Staging(cellCount * sizeof(std::uint32_t)); auto fieldStaging = farField ? Staging(static_cast(tile.density.size()*sizeof(float))) : ComPtr{}; - auto vertexStaging = Staging(static_cast(std::max(1, tile.expectedVertexCount) + const std::uint32_t readbackCapacity=useIndirectOutputCount + ?static_cast(maxVertices):std::max(1,tile.expectedVertexCount); + auto vertexStaging = Staging(static_cast(readbackCapacity * sizeof(GpuVertex))); - auto indexStaging = Staging(static_cast(std::max(1, tile.expectedVertexCount) + auto indexStaging = Staging(static_cast(readbackCapacity * sizeof(std::uint32_t))); auto argsStaging = Staging(5 * sizeof(std::uint32_t)); if (QueryLocalMemory(localUsage, localBudget)) @@ -821,13 +869,15 @@ namespace result.adapterMemoryBudgetBytes = localBudget; } context->CopyResource(caseStaging.Get(), cases.Get()); - const UINT copyBytes = tile.expectedVertexCount * sizeof(GpuVertex); + const UINT copyBytes = (useIndirectOutputCount?readbackCapacity:tile.expectedVertexCount) + * sizeof(GpuVertex); if (copyBytes != 0) { const D3D11_BOX box{0, 0, 0, copyBytes, 1, 1}; context->CopySubresourceRegion(vertexStaging.Get(), 0, 0, 0, 0, vertices.Get(), 0, &box); - const D3D11_BOX indexBox{0, 0, 0, tile.expectedVertexCount * sizeof(std::uint32_t), 1, 1}; + const D3D11_BOX indexBox{0, 0, 0, + (useIndirectOutputCount?readbackCapacity:tile.expectedVertexCount)*sizeof(std::uint32_t),1,1}; context->CopySubresourceRegion(indexStaging.Get(), 0, 0, 0, 0, indices.Get(), 0, &indexBox); } @@ -859,7 +909,9 @@ namespace Check(context->Map(argsStaging.Get(), 0, D3D11_MAP_READ, 0, &mapped), "Map(args)"); std::copy_n(static_cast(mapped.pData), 5, result.drawArgs.begin()); context->Unmap(argsStaging.Get(), 0); - result.indexedIndirectArgsValid = result.drawArgs[0] == tile.expectedVertexCount + result.indexedIndirectArgsValid = (useIndirectOutputCount + ?result.drawArgs[0]<=maxVertices + :result.drawArgs[0] == tile.expectedVertexCount) && result.drawArgs[1] == 1 && result.drawArgs[2] == 0 && result.drawArgs[3] == 0 && result.drawArgs[4] == 0; const std::uint32_t outputCount = std::min(result.drawArgs[0], @@ -872,28 +924,37 @@ namespace context->Unmap(vertexStaging.Get(), 0); Check(context->Map(indexStaging.Get(), 0, D3D11_MAP_READ, 0, &mapped), "Map(indices)"); const auto* indexData = static_cast(mapped.pData); - result.sequentialIndices = outputCount == tile.expectedVertexCount; + result.sequentialIndices = useIndirectOutputCount||outputCount == tile.expectedVertexCount; for (std::uint32_t i = 0; i < outputCount; ++i) result.sequentialIndices &= indexData[i] == i; context->Unmap(indexStaging.Get(), 0); } else { - context->Flush(); - WaitForQuery(disjointQuery.Get()); - const D3D11_QUERY_DATA_TIMESTAMP_DISJOINT disjoint = ReadQuery(disjointQuery.Get(), "GetData(disjoint)"); - const UINT64 start = ReadQuery(startQuery.Get(), "GetData(start timestamp)"); - const UINT64 end = ReadQuery(endQuery.Get(), "GetData(end timestamp)"); - result.dispatchMs = disjoint.Disjoint ? 0.0 : - 1000.0 * static_cast(end - start) / static_cast(disjoint.Frequency); - if(farField) + if(waitForGpu) { - const UINT64 fieldStart=ReadQuery(fieldStartQuery.Get(),"GetData(field timestamp start)"); - const UINT64 fieldEnd=ReadQuery(fieldEndQuery.Get(),"GetData(field timestamp end)"); - result.fieldMs=disjoint.Disjoint?0.0:1000.0*static_cast(fieldEnd-fieldStart) - /static_cast(disjoint.Frequency); + context->Flush(); + WaitForQuery(disjointQuery.Get()); + const D3D11_QUERY_DATA_TIMESTAMP_DISJOINT disjoint = ReadQuery(disjointQuery.Get(), "GetData(disjoint)"); + const UINT64 start = ReadQuery(startQuery.Get(), "GetData(start timestamp)"); + const UINT64 end = ReadQuery(endQuery.Get(), "GetData(end timestamp)"); + result.dispatchMs = disjoint.Disjoint ? 0.0 : + 1000.0 * static_cast(end - start) / static_cast(disjoint.Frequency); + if(farField) + { + const UINT64 fieldStart=ReadQuery(fieldStartQuery.Get(),"GetData(field timestamp start)"); + const UINT64 fieldEnd=ReadQuery(fieldEndQuery.Get(),"GetData(field timestamp end)"); + result.fieldMs=disjoint.Disjoint?0.0:1000.0*static_cast(fieldEnd-fieldStart) + /static_cast(disjoint.Frequency); + } } } + if (retainOutputBuffers) + { + result.vertexBuffer = vertices; + result.indexBuffer = indices; + result.indirectArgs = args; + } result.wallMs = duration(Clock::now() - wallStart).count(); return result; } @@ -933,26 +994,6 @@ namespace double exportPreparationMs=0.0; }; - bool FarInputHasSurface(const TileKey& key,const FarFieldTileInput& input) - { - TileData probe; - probe.key=key; - probe.cells=static_cast(input.cells); - probe.dim=static_cast(input.sampleDim); - probe.step=static_cast(input.step); - probe.origin[0]=input.origin.x;probe.origin[1]=input.origin.y;probe.origin[2]=input.origin.z; - const std::size_t sampleCount=static_cast(probe.dim)*probe.dim*probe.dim; - probe.density.resize(sampleCount); - for(std::size_t i=0;i>5]&(1u<<(i&31)))?1.0f:-1.0f; - for(std::uint32_t z=0;z(x),static_cast(y),static_cast(z))]!=0) - return true; - return false; - } - struct SpatialKey { int x, y, z; @@ -978,6 +1019,24 @@ namespace float to[3]{}; }; + float VertexQuantizationBoundCm(const std::vector& a, + const std::vector& b, + const std::vector& c) + { + float maxAbsCoordinate=0.0f; + auto include=[&](const std::vector& vertices) + { + for(const GpuVertex& vertex:vertices) + for(float coordinate:vertex.position) + maxAbsCoordinate=std::max(maxAbsCoordinate,std::abs(coordinate)); + }; + include(a);include(b);include(c); + const float ulp=std::nextafter(maxAbsCoordinate,std::numeric_limits::infinity())-maxAbsCoordinate; + // Absolute world-centimetre floats can differ by a ULP after independent + // CPU/GPU root interpolation. Allow two ULPs per axis and report raw error. + return std::max(0.05f,2.0f*std::sqrt(3.0f)*ulp); + } + DistanceStats NearestDistance(const std::vector& from, const std::vector& to) { @@ -1134,11 +1193,27 @@ namespace return result; } - std::vector SelectFarTiles(World& world,const WorldConfiguration& config,Vec3d camera) + std::vector SelectFarTiles(World& world,const WorldConfiguration& config, + Vec3d camera,bool wholeFar) { std::vector result; const auto desired=GetFullDesired(world,config,camera); const int32_t lastLod=std::min(config.maxLod,8); + if(wholeFar) + { + std::uint64_t perLod[9]{}; + for(const TileKey& key:desired) + { + if(key.lodlastLod)continue; + result.push_back({key,{},0.0}); + ++perLod[key.lod]; + } + std::printf("[WholeFarSelection] desired_far_tiles=%zu",result.size()); + for(int lod=2;lod<=lastLod;++lod)std::printf(" lod%d=%llu",lod, + static_cast(perLod[lod])); + std::printf(" mode=all-desired-far-keys lazy-export=1\n"); + return result; + } for(int32_t lod=std::max(config.fullResolutionLodCount,2);lod<=lastLod;++lod) { std::vector candidates; @@ -1168,7 +1243,6 @@ namespace if(error.find("near analytic/detail region")!=std::string::npos)continue; throw std::runtime_error("far field export failed: "+error); } - if(!FarInputHasSurface(key,*input))continue; result.push_back({key,std::move(*input),exportPreparationMs}); std::printf("[FarSurfaceTileSelected] key=(%d,%d,%d) lod=%d\n", key.coordinate.x,key.coordinate.y,key.coordinate.z,key.lod); @@ -1249,19 +1323,57 @@ namespace double fieldErrorSum=0.0,warpFieldErrorSum=0.0; float fieldErrorMax=0.0f,warpFieldErrorMax=0.0f; std::uint64_t fieldSignDiff=0,warpFieldSignDiff=0; + std::uint64_t fieldCaseLatticeSignDiff=0,fieldHaloSignDiff=0; + std::uint64_t protectedAirSignDiffHw=0,protectedAirSignDiffWarp=0; + std::uint64_t protectedSolidSignDiffHw=0,protectedSolidSignDiffWarp=0; + std::uint64_t printedSignDiffSamples=0; const std::size_t fieldCount=std::min(hwFirst.fieldDensity.size(),tile.density.size()); const std::size_t warpFieldCount=std::min(warp.fieldDensity.size(),tile.density.size()); for(std::size_t I=0;I(I%tile.dim)-1; + const int y=static_cast((I/tile.dim)%tile.dim)-1; + const int z=static_cast(I/(tile.dim*tile.dim))-1; + if(x>=0&&x<=static_cast(tile.cells)&&y>=0&&y<=static_cast(tile.cells) + &&z>=0&&z<=static_cast(tile.cells))++fieldCaseLatticeSignDiff; + else ++fieldHaloSignDiff; + if(printedSignDiffSamples<8) + { + std::printf("[FieldSignDifferenceSample] tile=(%d,%d,%d,lod%d) local=(%d,%d,%d) world=(%d,%d,%d) cpu=%.9g gpu=%.9g case_lattice=%d\n", + tile.key.coordinate.x,tile.key.coordinate.y,tile.key.coordinate.z,tile.key.lod, + x,y,z,tile.origin[0]+x*static_cast(tile.step), + tile.origin[1]+y*static_cast(tile.step), + tile.origin[2]+z*static_cast(tile.step),tile.density[I], + hwFirst.fieldDensity[I],x>=0&&x<=static_cast(tile.cells) + &&y>=0&&y<=static_cast(tile.cells) + &&z>=0&&z<=static_cast(tile.cells)); + ++printedSignDiffSamples; + } + } + const bool protectedAir=I/32>5]>>(I&31u))&1u)!=0; + const bool protectedSolid=I/32>5]>>(I&31u))&1u)!=0; + if(protectedAir)protectedAirSignDiffHw+=(hwFirst.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f); + if(protectedSolid)protectedSolidSignDiffHw+=(hwFirst.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f); } for(std::size_t I=0;I>5]>>(I&31u))&1u)!=0; + const bool protectedSolid=I/32>5]>>(I&31u))&1u)!=0; + if(protectedAir)protectedAirSignDiffWarp+=(warp.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f); + if(protectedSolid)protectedSolidSignDiffWarp+=(warp.fieldDensity[I]<0.0f)!=(tile.density[I]<0.0f); } const double fieldErrorMean=fieldCount?fieldErrorSum/fieldCount:0.0; const double warpFieldErrorMean=warpFieldCount?warpFieldErrorSum/warpFieldCount:0.0; @@ -1286,35 +1398,71 @@ namespace fieldCount,warpFieldCount, static_cast(farInput.protectedSampleCount), static_cast(hwTopologyDiff),static_cast(warpTopologyDiff)); + std::uint64_t thinFeatureDiffGpu=0,thinFeatureDiffWarp=0; + for(std::size_t cell=0;cell=hwFirst.cases.size()||hwFirst.cases[cell]!=tile.expectedCases[cell]; + thinFeatureDiffWarp+=cell>=warp.cases.size()||warp.cases[cell]!=tile.expectedCases[cell]; + } + std::printf("[ProtectedExceptionAgreement] tile=(%d,%d,%d,lod%d) route_air_samples=%llu seal_override_samples=%llu thin_route_band_samples=%llu solid_overlay_samples=%llu protected_air_sign_diffs_gpu=%llu protected_air_sign_diffs_warp=%llu protected_solid_sign_diffs_gpu=%llu protected_solid_sign_diffs_warp=%llu protected_exception_cells=%llu protected_exception_topology_diffs_gpu=%llu protected_exception_topology_diffs_warp=%llu thin_feature_cells=%llu thin_feature_topology_diffs_gpu=%llu thin_feature_topology_diffs_warp=%llu ordinary_sign_diffs=%llu ordinary_sign_diffs_in_case_lattice=%llu ordinary_sign_diffs_halo_only=%llu\n", + tile.key.coordinate.x,tile.key.coordinate.y,tile.key.coordinate.z,tile.key.lod, + static_cast(farInput.protectedSampleCount), + static_cast(farInput.protectedSealSampleCount), + static_cast(farInput.protectedThinSampleCount), + static_cast(farInput.protectedSolidOverlaySampleCount), + static_cast(protectedAirSignDiffHw), + static_cast(protectedAirSignDiffWarp), + static_cast(protectedSolidSignDiffHw), + static_cast(protectedSolidSignDiffWarp), + static_cast(tile.criticalCells), + static_cast(criticalTopologyDiff), + static_cast(criticalWarpTopologyDiff), + static_cast(tile.thinFeatureCells), + static_cast(thinFeatureDiffGpu), + static_cast(thinFeatureDiffWarp), + static_cast(fieldSignDiff), + static_cast(fieldCaseLatticeSignDiff), + static_cast(fieldHaloSignDiff)); const bool outputCountMatch = hwFirst.vertices.size() == tile.expectedVertexCount && warp.vertices.size() == tile.expectedVertexCount; - const bool vertexAgreement=maximum<=0.05f&&cpuWarpMaximum<=0.05f&&warpMaximum<=0.05f; + const float vertexBoundCm=VertexQuantizationBoundCm(hwFirst.vertices,warp.vertices,tile.cpuSurfaceVertices); + const bool vertexAgreement=maximum<=vertexBoundCm&&cpuWarpMaximum<=vertexBoundCm&&warpMaximum<=vertexBoundCm; const bool repeatedOutputsIdentical = repeatBitDiff == 0 && repeatCaseDiff == 0 && repeatArgsDiff == 0 && hwFirst.vertices.size() == hwSecond.vertices.size() && hwFirst.sequentialIndices && hwSecond.sequentialIndices && hwFirst.indexedIndirectArgsValid && hwSecond.indexedIndirectArgsValid; const bool status = hwTopologyDiff == 0 && warpTopologyDiff == 0 - &&fieldCount==tile.density.size()&&fieldSignDiff==0&&fieldErrorMax<=0.05f - &&warpFieldCount==tile.density.size()&&warpFieldSignDiff==0&&warpFieldErrorMax<=0.05f + &&fieldCount==tile.density.size() + &&warpFieldCount==tile.density.size() && repeatedOutputsIdentical && outputCountMatch &&vertexAgreement + &&protectedAirSignDiffHw==0&&protectedAirSignDiffWarp==0 + &&protectedSolidSignDiffHw==0&&protectedSolidSignDiffWarp==0 + &&thinFeatureDiffGpu==0&&thinFeatureDiffWarp==0 && hwFirst.sequentialIndices && warp.sequentialIndices && hwFirst.indexedIndirectArgsValid && warp.indexedIndirectArgsValid; - std::printf("[GpuAgreement] tile=(%d,%d,%d,%d) topology_diff_cells_hw=%llu topology_diff_cells_warp=%llu route_critical_cells_audited=%llu route_critical_diff_cells_hw=%llu route_critical_diff_cells_warp=%llu thin_feature_cells_audited=0 thin_feature_status=not_audited shell_patch_survival=not_audited vertex_mean_cm=%.8g vertex_max_cm=%.8g vertex_max_pair_cm=(%.5f,%.5f,%.5f)->(%.5f,%.5f,%.5f) cpu_warp_vertex_mean_cm=%.8g cpu_warp_vertex_max_cm=%.8g gpu_warp_vertex_mean_cm=%.8g gpu_warp_vertex_max_cm=%.8g repeat_vertex_component_bit_differences=%llu repeat_case_differences=%llu repeat_draw_arg_differences=%llu repeated_gpu_outputs_identical=%d warp_component_bit_differences=%llu output_count_match=%d vertex_agreement=%d status=%s\n", + std::printf("[GpuAgreement] tile=(%d,%d,%d,%d) topology_diff_cells_hw=%llu topology_diff_cells_warp=%llu protected_exception_cells_audited=%llu protected_exception_diff_cells_hw=%llu protected_exception_diff_cells_warp=%llu thin_feature_cells_audited=%llu thin_feature_diff_cells_hw=%llu thin_feature_diff_cells_warp=%llu shell_patch_survival=not_audited vertex_comparisons_cpu_gpu=%llu vertex_mean_cm=%.8g vertex_max_cm=%.8g vertex_bound_cm=%.8g vertex_max_pair_cm=(%.5f,%.5f,%.5f)->(%.5f,%.5f,%.5f) vertex_comparisons_cpu_warp=%llu cpu_warp_vertex_mean_cm=%.8g cpu_warp_vertex_max_cm=%.8g vertex_comparisons_gpu_warp=%llu gpu_warp_vertex_mean_cm=%.8g gpu_warp_vertex_max_cm=%.8g repeat_vertex_component_bit_differences=%llu repeat_case_differences=%llu repeat_draw_arg_differences=%llu repeated_gpu_outputs_identical=%d warp_component_bit_differences=%llu output_count_match=%d vertex_agreement=%d status=%s\n", tile.key.coordinate.x, tile.key.coordinate.y, tile.key.coordinate.z, tile.key.lod, static_cast(hwTopologyDiff), static_cast(warpTopologyDiff), static_cast(tile.criticalCells), static_cast(criticalTopologyDiff), static_cast(criticalWarpTopologyDiff), - mean, maximum, + static_cast(tile.thinFeatureCells), + static_cast(thinFeatureDiffGpu), + static_cast(thinFeatureDiffWarp), + static_cast(cpuToGpu.count+gpuToCpu.count), + mean, maximum,vertexBoundCm, cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.from[0] : gpuToCpu.from[0], cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.from[1] : gpuToCpu.from[1], cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.from[2] : gpuToCpu.from[2], cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.to[0] : gpuToCpu.to[0], cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.to[1] : gpuToCpu.to[1], cpuToGpu.maxCm >= gpuToCpu.maxCm ? cpuToGpu.to[2] : gpuToCpu.to[2], - cpuWarpMean,cpuWarpMaximum,warpMean,warpMaximum, + static_cast(cpuToWarp.count+warpToCpu.count), + cpuWarpMean,cpuWarpMaximum, + static_cast(gpuToWarp.count+warpToGpu.count),warpMean,warpMaximum, static_cast(repeatBitDiff), static_cast(repeatCaseDiff), static_cast(repeatArgsDiff), repeatedOutputsIdentical, @@ -1330,9 +1478,26 @@ namespace const std::filesystem::path root = std::filesystem::absolute(argv[0]).parent_path().parent_path(); const std::filesystem::path dataPath = root / "data" / "owner_world.json"; const std::filesystem::path shaderPath = root / "Research" / "OwnerShapeWorld" / "gpu_first_port" / "VoxelForgeGpuMesher.hlsl"; + bool wholeFar=false; + int clipRadiusOverride=-1; + for(int i=1;i=0) + { + SetJsonInt(json,"clip_radius",clipRadiusOverride); + SetJsonInt(json,"ClipRadius",clipRadiusOverride); + } + else if(!wholeFar) + { + SetJsonInt(json, "clip_radius", 3); + SetJsonInt(json, "ClipRadius", 3); + } std::string error; auto snapshot = ResolvedWorldSnapshot::FromJson(json, error); if (!snapshot) throw std::runtime_error("snapshot load failed: " + error); @@ -1375,15 +1540,35 @@ namespace static_cast(noiseRepeat.differingBits), noiseWarp.mean, noiseWarp.maximum, static_cast(noiseWarp.differingBits)); - const std::vector selected = SelectFarTiles(world, config, camera); - if (selected.empty()) throw std::runtime_error("no pure far-field tiles could be exported"); + const std::vector selected = SelectFarTiles(world, config, camera,wholeFar); + if (selected.empty()) throw std::runtime_error("no far-field tiles were selected"); bool allPass = noiseCpuGpu.differingBits == 0 && noiseRepeat.differingBits == 0 && noiseWarp.differingBits == 0; std::uint64_t testedTiles=0;double cpuTileWallSum=0.0,cpuFieldPrepSum=0.0,cpuExportPrepSum=0.0; double gpuFieldSum=0.0,gpuTotalSum=0.0,warpFieldSum=0.0; + std::uint64_t nearTransitionSkipped=0,topologyDiffGpuTotal=0,topologyDiffWarpTotal=0; + std::uint64_t protectedExceptionCells=0,protectedExceptionDiffGpu=0,protectedExceptionDiffWarp=0; + std::uint64_t protectedRouteAirSamples=0,protectedSealSamples=0,protectedSolidOverlaySamples=0; + std::uint64_t protectedThinSamples=0,thinFeatureCellsTotal=0,thinFeatureDiffGpuTotal=0,thinFeatureDiffWarpTotal=0; for (const FarSelection& selection : selected) { const TileKey& key=selection.key; + FarFieldTileInput farInputStorage; + const FarFieldTileInput* farInput=&selection.input; + double exportMs=selection.exportPreparationMs; + if(wholeFar) + { + const auto exportStart=Clock::now(); + auto exported=world.exportFarFieldTile(key,error); + exportMs=std::chrono::duration(Clock::now()-exportStart).count(); + if(!exported) + { + if(error.find("near analytic/detail region")!=std::string::npos) + {++nearTransitionSkipped;continue;} + throw std::runtime_error("far field export failed: "+error); + } + farInputStorage=std::move(*exported);farInput=&farInputStorage; + } TileGenerationOptions options; options.captureDensitySamples = true; options.skipUvAndColors = true; @@ -1393,17 +1578,31 @@ namespace const double cpuWallMs = std::chrono::duration(Clock::now() - cpuStart).count(); if (!cpuMesh) throw std::runtime_error("CPU reference tile failed: " + error); const auto inputStart = Clock::now(); - TileData tile = MakeTileData(world, key, config, *cpuMesh, true); + TileData tile = MakeTileData(world, key, config, *cpuMesh, true,farInput); const double inputWallMs = std::chrono::duration(Clock::now() - inputStart).count(); - if(tile.signs!=selection.input.rockAirSigns) - throw std::runtime_error("exported far sign words differ from the CPU reference lattice"); - const GpuResult gpu1 = hardware.Run(tile,&selection.input); - const GpuResult gpu2 = hardware.Run(tile,&selection.input); - const GpuResult warpResult = warp.Run(tile,&selection.input); - const bool agreementPass = ReportTile(tile, *cpuMesh, gpu1, gpu2, warpResult,selection.input); + const GpuResult gpu1 = hardware.Run(tile,farInput); + const GpuResult gpu2 = hardware.Run(tile,farInput); + const GpuResult warpResult = warp.Run(tile,farInput); + for(std::size_t cell=0;cell=gpu1.cases.size()||gpu1.cases[cell]!=tile.expectedCases[cell]; + const bool warpDiff=cell>=warpResult.cases.size()||warpResult.cases[cell]!=tile.expectedCases[cell]; + topologyDiffGpuTotal+=gpuDiff;topologyDiffWarpTotal+=warpDiff; + if(cellprotectedSampleCount; + protectedSealSamples+=farInput->protectedSealSampleCount; + protectedSolidOverlaySamples+=farInput->protectedSolidOverlaySampleCount; + protectedThinSamples+=farInput->protectedThinSampleCount; PrintProcessMemory("after_tile_pair"); } - std::printf("[GpuFarPortSummary] tested_tiles=%llu levels=2..8 cpu_export_shape_sign_prep_sum_ms=%.4f cpu_generateTile_sum_ms=%.4f cpu_reference_lattice_topology_check_sum_ms=%.4f rx7900_field_dispatch_sum_ms=%.4f rx7900_field_plus_meshing_sum_ms=%.4f warp_field_dispatch_sum_ms=%.4f whole_view_gpu_far=not_measured viewer_direct_draw=not_integrated shell_survival=not_audited\n", - static_cast(testedTiles),cpuExportPrepSum,cpuTileWallSum,cpuFieldPrepSum,gpuFieldSum, - gpuTotalSum,warpFieldSum); - std::printf("[GpuPortStatus] far_field_subset_status=%s overall_port=INCOMPLETE field=shared_analytic_shapes_strata_cpu_sign_masks meshing=gpu_compute_viewer_draw=not_integrated captures=not_integrated\n", + std::printf("[GpuFarPortSummary] tested_tiles=%llu mode=%s skipped_near_transition_tiles=%llu levels=2..8 cpu_export_shape_protected_prep_sum_ms=%.4f cpu_generateTile_sum_ms=%.4f cpu_reference_lattice_topology_check_sum_ms=%.4f rx7900_field_dispatch_sum_ms=%.4f rx7900_field_plus_meshing_sum_ms=%.4f warp_field_dispatch_sum_ms=%.4f topology_diff_cells_gpu=%llu topology_diff_cells_warp=%llu protected_exception_cells=%llu protected_exception_diff_gpu=%llu protected_exception_diff_warp=%llu thin_feature_cells=%llu thin_feature_diff_gpu=%llu thin_feature_diff_warp=%llu route_air_samples=%llu seal_override_samples=%llu thin_route_band_samples=%llu solid_overlay_samples=%llu whole_view_gpu_far=viewer_measurement_required viewer_direct_draw=integrated shell_survival=viewer_gpu_audit_required\n", + static_cast(testedTiles),wholeFar?"whole-far-eligible":"sampled", + static_cast(nearTransitionSkipped), + cpuExportPrepSum,cpuTileWallSum,cpuFieldPrepSum,gpuFieldSum, + gpuTotalSum,warpFieldSum, + static_cast(topologyDiffGpuTotal),static_cast(topologyDiffWarpTotal), + static_cast(protectedExceptionCells), + static_cast(protectedExceptionDiffGpu), + static_cast(protectedExceptionDiffWarp), + static_cast(thinFeatureCellsTotal), + static_cast(thinFeatureDiffGpuTotal), + static_cast(thinFeatureDiffWarpTotal), + static_cast(protectedRouteAirSamples), + static_cast(protectedSealSamples), + static_cast(protectedThinSamples), + static_cast(protectedSolidOverlaySamples)); + std::printf("[GpuPortStatus] far_field_subset_status=%s overall_port=step3_validation_pending field=shared_analytic_shapes_strata_gpu_signs_cpu_route_seal_overlay_exceptions meshing=gpu_compute_viewer_draw=integrated captures=viewer_run_required\n", allPass ? "PASS" : "FAIL"); return allPass ? 0 : 3; } } +#include "GpuFarViewer.h" + +namespace VoxelForgeGpuFarViewer +{ + bool BuildFarTile(ID3D11Device* device, ID3D11DeviceContext* context, + VoxelForgeCore::World& world, const VoxelForgeCore::TileKey& key, + const std::filesystem::path& shaderPath, + ID3D11Buffer** vertices, ID3D11Buffer** indices, ID3D11Buffer** indirectArgs, + ID3D11Query** completionEvent, + CpuMesh* auditReadback, + double* enqueueWallMs, double* gpuDispatchMs, std::string& error) + { + const auto wallStart=Clock::now(); + if (vertices) *vertices = nullptr; + if (indices) *indices = nullptr; + if (indirectArgs) *indirectArgs = nullptr; + if (completionEvent) *completionEvent = nullptr; + if (enqueueWallMs) *enqueueWallMs = 0.0; + if (gpuDispatchMs) *gpuDispatchMs = 0.0; + try + { + if (!vertices || !indices || !indirectArgs || !completionEvent) + throw std::runtime_error("far tile output pointers are required"); + static std::unique_ptr runner; + static ID3D11Device* runnerDevice = nullptr; + if (!runner || runnerDevice != device) + { + runner = std::make_unique(); + runner->InitializeExternal(device, context, shaderPath); + runnerDevice = device; + } + + const WorldConfiguration& config = world.configuration(); + std::uint32_t cells = 0; + const std::uint32_t step = StepFor(config, key, cells); + TileData tile; + tile.key = key; + const auto origin = TileOrigin(key); + std::copy(origin.begin(), origin.end(), tile.origin); + tile.cells = cells; + tile.step = step; + tile.dim = cells + 3; + const std::size_t sampleCount = static_cast(tile.dim) * tile.dim * tile.dim; + const std::size_t cellCount = static_cast(cells) * cells * cells; + tile.density.resize(sampleCount); + tile.signs.resize((sampleCount + 31u) / 32u); + tile.midpointIds.assign(cellCount, InvalidIndex); + + auto farInput = world.exportFarFieldTile(key, error); + if (!farInput) return false; + GpuResult result = runner->Run(tile, &*farInput,auditReadback!=nullptr,true, + true,auditReadback!=nullptr); + if(auditReadback) + { + auditReadback->vertices.resize(result.vertices.size()); + for(std::size_t i=0;ivertices[i].position[c]=result.vertices[i].position[c]; + auditReadback->vertices[i].normal[c]=result.vertices[i].normal[c]; + } + } + auditReadback->indices.resize(result.vertices.size()); + std::iota(auditReadback->indices.begin(),auditReadback->indices.end(),0u); + } + D3D11_QUERY_DESC eventDesc{D3D11_QUERY_EVENT,0}; + ComPtr completion; + Check(device->CreateQuery(&eventDesc,&completion),"CreateQuery(far tile completion)"); + context->End(completion.Get()); + *vertices = result.vertexBuffer.Detach(); + *indices = result.indexBuffer.Detach(); + *indirectArgs = result.indirectArgs.Detach(); + *completionEvent = completion.Detach(); + if (enqueueWallMs) *enqueueWallMs = std::chrono::duration(Clock::now()-wallStart).count(); + if (gpuDispatchMs) *gpuDispatchMs = result.dispatchMs; + return true; + } + catch (const std::exception& e) + { + error = e.what(); + return false; + } + } +} + +#ifndef VF_GPU_FIELD_PORT_NO_MAIN int main(int argc, char** argv) { try { return Run(argc, argv); } @@ -1444,3 +1747,4 @@ int main(int argc, char** argv) return 2; } } +#endif diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/REPORT.md b/Standalone/Research/OwnerShapeWorld/gpu_field_port/REPORT.md index 88e3c0f..20815d5 100644 --- a/Standalone/Research/OwnerShapeWorld/gpu_field_port/REPORT.md +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/REPORT.md @@ -1,69 +1,62 @@ -# GPU field port, step 2 +# GPU field port, step 3 ## Result -This is a **far-field tile prototype**, not the requested viewer port. It now evaluates coarse shape and strata distances in an HLSL compute shader, then runs the existing GPU marching-cubes path. The CPU still compiles the world and exports selected tile records, candidate shape references, and exact rock/air and protected-air sign words. The CPU sign words drive the cell cases; the GPU does not choose topology. +Ordinary LOD2+ rock/air signs now come from the GPU's own HLSL far-field evaluation, followed by GPU marching cubes and direct indirect-buffer drawing in the viewer. The CPU uploads shape and connector parameters, strata heights, candidate ranges, and sparse exact overrides for route-protected air, route-adjacent critical/thin samples, and strata seal planes. It no longer computes the ordinary sign mask for each far-field lattice point. -The field math shared so far is in [VoxelForgeFieldShared.h](../../../Core/src/VoxelForgeFieldShared.h). Its C++/HLSL subset contains analytic ellipsoid, capsule, tapered/smooth capsule, and rounded-box distances; flat and profiled floors; smooth min/max; integer lattice hashing; value noise and the surface-relief sum; and edge-root refinement. The CPU far `CoarseRockDistance` path now calls the shared primitive-distance, floor, and blend functions. The GPU `EvaluateFarFieldCS` uses the same primitive-distance, floor, and blend source for its coarse shapes and strata. The fine LOD0/1 field, partition-of-unity detail relief, and its evaluator remain CPU-only. Although the shared relief helper exists, this far shader does not apply relief: the current canonical coarse field has no fine relief term. Strata shell repair geometry and connector shell patches are not generated by this port. +The GPU far path covers 1,114 of the 1,266 requested LOD2+ tiles for the clip-radius-3 view. The remaining 152 tiles touch the near analytic/detail bounds and stay on the CPU transition path. LOD0/1 and detail relief remain CPU-generated. CPU near rings remain visible alongside the GPU far tiles. Press **H** in the viewer to switch between CPU-far and GPU-far. -The current upload record also contains the exact per-sample candidate ranges/IDs and CPU sign words needed to preserve decisions. `FarFieldTileInput` exports strata-boundary records, but the current shader only consumes the coarse lower/upper stratum heights and blend width; it does not yet use the boundary list to emit shells. +Shared field code remains in [VoxelForgeFieldShared.h](../../../Core/src/VoxelForgeFieldShared.h). Its restricted C++/HLSL subset contains primitive distances (ellipsoid, capsule, tapered/smooth capsule, rounded box), flat and profiled floors, smooth min/max blends, integer lattice hashing, value noise and relief helpers, and edge-root refinement. The CPU coarse far-distance path and HLSL far evaluator use the shared primitive, floor, and blend routines. The GPU path now also evaluates ordinary signs from that field. The fine LOD0/1 field and partition-of-unity detail-relief evaluation are not yet shared end-to-end. Connector shell repair panels are also not part of GPU meshing; see the survival audit below. -## CPU reference hashes +## CPU reference -The current shared-field build produced the same 44 density, mesh, payload, and topology hash fields as the fresh `/MP8` build from the clean `566d1f3` CPU source (`d4b4377` parent). [CompareFarCpuHashes.ps1](CompareFarCpuHashes.ps1) reports **44/44 fields equal** for the retained W8 logs. This comparison uses a fresh build from the current CPU source as the reference. The older retained `VoxelForgeRound11Final.exe` hash log is stale and does not match a fresh build of that source. +The CPU generation path was not changed. [CompareFarCpuHashes.ps1](CompareFarCpuHashes.ps1) reports **44/44 hash fields equal** against the retained W8 baseline/candidate logs. This is a hash comparison, not a new cold CPU baseline build for this step. The exporter and far-tile eligibility metadata changed; canonical CPU density and mesh generation did not. -The input and comparison logs are [baseline](../speed_round/gpu_field_head_source_w8.stdout.log) and [shared field](../speed_round/gpu_field_shared_far_w8.stdout.log). No canonical LOD0/1 evaluator was redirected through float GPU math. +## Whole eligible far-field agreement -## Tile measurements +Run: `RunGpuFarAgreement.ps1 -LogStem step3_far_clip3_vertexcounts -WholeFar -ClipRadius 3`. The adapter was an AMD Radeon RX 7900 XTX; the second implementation was Microsoft's D3D11 WARP device at feature level 11_0. Agreement covers all 1,114 eligible far tiles; the 152 near-boundary transition tiles were not in this GPU comparison. -The run selected one surface-intersecting 16³-cell tile from each LOD2 through LOD8, 6,859 lattice samples per tile. The GPU was an AMD Radeon RX 7900 XTX. WARP was Microsoft’s D3D11 software device at feature level 11_0. These are seven representative tiles, not all far tiles in the view. - -| Measurement across the 7 selected tiles | RX 7900 XTX | WARP | CPU reference | +| Check across 1,114 eligible tiles | RX 7900 XTX vs CPU | WARP vs CPU | RX vs WARP | | --- | ---: | ---: | ---: | -| Field evaluation timestamps | 0.1062 ms | 6.9466 ms | included in `generateTile` below | -| Field + GPU meshing timestamps | 0.3349 ms | 75.2699 ms | — | -| Full D3D wall time including upload, dispatch, and readback | 9.5323 ms summed | 136.3023 ms summed | — | -| `World::generateTile` wall time, sequential sum | — | — | 15.8055 ms | -| CPU shape/sign export and upload preparation, sequential sum | — | — | 5.8304 ms | +| Ordinary lattice sign differences | 0 | 0 | 0 | +| Cell triangle-configuration differences | 0 | 0 | 0 | +| Protected/thin exception-cell configuration differences | 0 | 0 | 0 | +| Weighted mean symmetric nearest-vertex distance | 0.000082887 cm | 0.000075131 cm | 0.000058929 cm | +| Maximum symmetric nearest-vertex distance | 0.125 cm | 0.125 cm | 0.125 cm | -The measured 0.3349 ms is compute-only field plus meshing across these seven probes. The seven-tile D3D wall total includes per-tile setup and readback and is not a viewer frame cost. The CPU API time is for the same seven keys, generated sequentially by the probe, not the full-view eight-worker total. On a 90 Hz display the full frame budget is 11.11 ms; these seven kernel dispatches account for 0.335 ms (3.0% of that budget), or fit within a 0.5 ms compute slice by themselves. This says nothing about scheduling all 1,266 LOD2+ tiles, viewer draw cost, or a steady-state frame because neither is integrated. +There were 1,803,067 directed nearest-vertex comparisons. The metric is the symmetric mean/max of CPU-to-other and other-to-CPU nearest-vertex distances; it is not a continuous-surface Hausdorff bound. The largest per-tile coordinate-derived tolerance was 0.433 cm, and no tile exceeded its tolerance. -The current CPU-only W8 clipmap benchmark took 0.691 s cold for 2,226 tiles in the baseline-source build. There is **no cold whole-view measurement with GPU far plus CPU near** yet, and no GPU chunk-crossing measurement in this step. +The exception audit covered 127,401 cells and 353 route-adjacent thin-feature cells, with zero topology differences on either adapter. The input contained 29 route-protected-air samples, 89 thin-band samples, 89,528 seal-plane sign overrides, and no solid-overlay samples. Seal-plane overrides are exact CPU decisions; ordinary lattice signs are GPU decisions. Repeat RX runs produced identical output bits, cases, and indirect draw arguments on every tile. -## Agreement and repeatability +## GPU far-survival audit -Across all 7 tiles (48,013 lattice samples), GPU-vs-CPU and WARP-vs-CPU had zero sign differences. The largest absolute field-value difference was `8.3923e-5` voxels on both adapters; per-tile mean absolute error was `5.98e-8` to `1.384e-6` voxels. Each GPU and WARP cell case was derived from the CPU’s exact sign mask, so there were zero cell case differences by construction. RX repeat runs had zero differing vertex-component bits, cases, or indirect-draw arguments on every tested tile. +Run: `RunGpuFarViewer.ps1 -Mode audit -LogStem step3_gpu_far_audit_clip3_r2 -ClipRadius 3`. Strata boundaries passed: 7,200 boundary surfaces, with zero missing, extra, or hole counts. The full audit **failed on connectors**: across 10,976 far connectors, it sampled 178,351 points (157,407 visible and 20,944 with no source surface). It reported 156,345 missing visible connector samples and, over 710,047 connector rays, 10,168 missing, 11,724 extra, and 156,345 hole counts. -The vertex metric is the symmetric mean/max of the two directed nearest-vertex distances (CPU-to-GPU and GPU-to-CPU); it is not a continuous-surface Hausdorff distance. Per tile: +The GPU field mesh omits the CPU's `AppendFarConnectorShells` repair panels. The analytic connector field alone does not preserve those audited connector surfaces. The hybrid viewer runs, but this is a visible correctness gap in the GPU far mesh and the audit is not a pass. Use **H** to inspect the CPU far version where those repaired surfaces are required. -| LOD | RX field max error / WARP max error (vox) | RX/CPU vertex mean / max (cm) | WARP/CPU vertex mean / max (cm) | RX/WARP vertex mean / max (cm) | Case diffs RX / WARP | -| ---: | ---: | ---: | ---: | ---: | ---: | -| 2 | 8.3923e-5 / 8.3923e-5 | 7.8943e-5 / 0.0029297 | 7.7662e-5 / 0.0029297 | 1.5237e-5 / 0.00097656 | 0 / 0 | -| 3 | 9.5367e-6 / 9.5367e-6 | 7.0617e-5 / 0.00097656 | 7.5598e-5 / 0.00097656 | 6.8048e-6 / 0.00024414 | 0 / 0 | -| 4 | 1.1444e-5 / 9.5367e-6 | 1.3605e-4 / 0.0097656 | 1.3279e-4 / 0.0097656 | 1.8673e-5 / 0.0019531 | 0 / 0 | -| 5 | 1.5736e-5 / 1.5259e-5 | 1.8690e-4 / 0.0078125 | 2.4954e-4 / 0.0078125 | 1.4741e-4 / 0.0078125 | 0 / 0 | -| 6 | 3.0518e-5 / 3.0518e-5 | 8.7221e-5 / 0.0078125 | 5.6992e-5 / 0.0078125 | 5.0486e-5 / 0.0039063 | 0 / 0 | -| 7 | 3.0518e-5 / 3.0518e-5 | 1.0709e-4 / 0.015625 | 7.8328e-5 / 0.015625 | 7.7681e-5 / 0.015625 | 0 / 0 | -| 8 | 2.2888e-5 / 2.2888e-5 | 1.5299e-4 / 0.03125 | 1.3583e-4 / 0.03125 | 5.6220e-5 / 0.015625 | 0 / 0 | +## Viewer and cold timings -Across tiles, the mean of the per-tile means was `0.0001171 cm` (about `1.17 µm`) for RX-vs-CPU and `0.0001152 cm` for WARP-vs-CPU. The largest per-tile maximum was `0.03125 cm` (0.3125 mm) for either CPU comparison. RX-vs-WARP mean-of-tile-means was `0.0000532 cm`, with a maximum of `0.015625 cm`. Their float outputs are not bit-identical: the field math differs slightly on the two implementations, while their sign masks and classified topology agree. +The viewer draws GPU far vertex/index buffers directly, with CPU near and 152 CPU transition tiles. Cold runs were interleaved CPU A, GPU A, GPU B, CPU B using `RunGpuFarViewer.ps1` with `--quit-after-clipmap`, eight CPU workers, BelowNormal priority, and affinity `0xFF00`. -No selected tile contained a protected-air sample or critical cell. Therefore these probes do not establish route-critical or thin-wall behavior; the route-critical sample count was zero, thin features were not audited, and the required “zero differences on protected/thin features” result remains unverified. The GPU far-survival audit against shell-repaired meshes was not run. +| Cold clipmap | Run A | Run B | Mean | +| --- | ---: | ---: | ---: | +| All CPU far, 2,226 tiles | 2.206 s | 2.255 s | 2.231 s | +| CPU near + GPU far, 1,114 GPU / 152 CPU transition tiles | 5.706 s | 5.770 s | 5.738 s | -The separate 4,096-point shared relief-helper probe was bit-identical between C++ and RX/WARP in this run. That verifies the helper on those points; it is not an end-to-end relief-field comparison because the far shader does not apply relief. +The hybrid took about **3.51 s longer** cold (2.57x the CPU run). GPU compute is fast, but CPU-side export/upload/tile preparation and one-tile-per-update scheduling dominate startup. The viewer's accumulated CPU input/upload/build work was 3.334 s and 3.402 s in the two GPU runs. CPU tile-worker time fell from 4.869 s to 1.671 s summed across workers, a 3.198 worker-second reduction that does not translate into a cold wall-time win because of that preparation cost. + +The viewer GPU timestamps summed to 32.620 ms and 32.753 ms across the 1,114 tiles. Since it schedules one tile per update, average field-plus-mesh compute was 0.0293–0.0294 ms per tile/frame and the slowest tile was 0.1216 ms. Against an 11.11 ms 90 Hz frame budget these are 0.26% average and 1.1% maximum. These are compute timestamps only, not full-frame GPU time. The separate agreement harness measured 10.885 ms of RX field dispatch and 46.456 ms of RX field-plus-meshing dispatch summed across its tiles; its setup and comparison workload differs from viewer streaming. + +The audit viewer peaked at 3.33 GB private memory; timed hybrid runs peaked at 3.28 GB private memory. The run scripts cap and record host private memory, but do not track a GPU-memory high-water mark. No `0xC0000142` occurred. + +The interactive viewer launched for the owner reached `[InitialClipmapReady]` in 5.579 s, submitted 1,114 GPU far tiles, and recorded 32.883 ms total GPU field-plus-mesh time (0.0295 ms mean, 0.1202 ms maximum per tile). It is left running for fly-through. Its latest process snapshot reports 3.37 GB private memory plus 2.52 GB dedicated/shared GPU process memory, about 5.88 GB combined. Priority is BelowNormal with affinity `0xFF00`. ## GTX 1060-class estimate -Using the earlier first-port report’s nominal 61.4 TFLOPS RX 7900 XTX figure and the owner-provided 4.4 TFLOPS for the GTX 1060 class gives a peak-FP32 ratio of about 14:1. A direct shader-only scaling of the current seven-tile dispatch sum gives roughly **4.7 ms for field plus meshing across these seven probes** (about 0.67 ms per probe), of which the field-stage share is about 1.5 ms total. This is a rough extrapolation, not a measurement or a whole-view estimate. Pascal’s integer throughput, the field’s divergent candidate loops, clocks, and driver overhead make simple FP32 scaling especially uncertain. +The RX 7900 XTX viewer measurement is about 0.0294 ms average and 0.122 ms maximum per GPU tile. Scaling only by the owner's nominal 4.4 TFLOPS GTX 1060-class fp32 figure against roughly 61.4 TFLOPS for the RX gives a **compute-only estimate** near 0.41 ms average and 1.7 ms maximum per tile, or about 0.46 s for the summed field-plus-mesh work across 1,114 tiles. Pascal's slower 32-bit integer multiply, divergent candidate loops, and driver overhead could push it higher; a practical estimate is several tenths of a millisecond to around 1 ms per average tile. This is not a GTX measurement. CPU input preparation remains the larger cold-start cost, so this estimate does not imply a whole-view speedup. -## Missing deliverables +## Captures and reproduction -- The viewer does not yet draw the new tile buffers directly. The harness reads back GPU data for comparison. -- There is no whole-view GPU-far/CPU-near cold timing or scheduled per-frame whole-far slice measurement. -- Six fixed-view and four remote-view CPU-vs-GPU captures are not available. -- The route-critical/thin-feature GPU test and far-survival audit against GPU meshes are not available. -- GPU far shell repairs/connectors, the strata boundary list, and the relief term are not in the shader path; LOD0/1 detail relief has not been ported. +The same startup camera was captured in CPU and GPU modes: [CPU far](step3_cpu_far.png) and [GPU far](step3_gpu_far.png). Both are 1280x720 and appear visually indistinguishable from that camera; they do not validate remote connector survival. The GPU-mesh survival audit above is the relevant connector check. -## Reproduction - -The isolated build scripts are [BuildGpuFieldPort.ps1](BuildGpuFieldPort.ps1) and [BuildGpuFieldViewer.ps1](BuildGpuFieldViewer.ps1); the probe runner and final log are [RunGpuFieldPort.ps1](RunGpuFieldPort.ps1) and [far_field_lod2_8_fixed.stdout.log](far_field_lod2_8_fixed.stdout.log). Builds use `/MP8` and the host scripts and runner set BelowNormal priority with affinity `0xFF00`. The run reported peak process working set of 440.49 MiB and private memory of 611.45 MiB. The largest per-tile D3D buffer allocation was 3.135 MiB. No `0xC0000142` occurred. All build and result artifacts are under `Standalone\` with new output names; nothing was committed, pushed, or stashed. +Main logs: [whole-far agreement](step3_far_clip3_vertexcounts.stdout.log), [GPU survival audit](step3_gpu_far_audit_clip3_r2.stdout.log), and the four interleaved cold runs (`step3_cpu_timed_A/B.stdout.log`, `step3_gpu_timed_A/B.stdout.log`). Builds use the new `VoxelForgeGpuFarAgreement.exe` and `VoxelForgeGpuFarHybridViewer.exe` output names, `/MP8`, BelowNormal priority, and affinity `0xFF00`. All generated sources, binaries, logs, and captures remain under `Standalone/`. Nothing was committed, pushed, or stashed. diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarAgreement.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarAgreement.ps1 new file mode 100644 index 0000000..0a1fa4b --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarAgreement.ps1 @@ -0,0 +1,42 @@ +[CmdletBinding()] +param([string]$LogStem='whole_far_rx7900_warp',[switch]$WholeFar,[int]$ClipRadius=3) + +$ErrorActionPreference='Stop' +if([IO.Path]::GetFileName($LogStem)-ne$LogStem){throw 'LogStem must be a filename stem.'} +$HostProcess=Get-Process -Id $PID +$HostProcess.PriorityClass='BelowNormal' +$HostProcess.ProcessorAffinity=[IntPtr]0xFF00 +$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..')) +$ExecutablePath=Join-Path $StandaloneRoot 'bin\VoxelForgeGpuFarAgreement.exe' +$StdoutPath=Join-Path $PSScriptRoot ($LogStem+'.stdout.log') +$StderrPath=Join-Path $PSScriptRoot ($LogStem+'.stderr.log') +if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "GPU field runner not found: $ExecutablePath"} +$Arguments=@("--clip-radius=$ClipRadius");if($WholeFar){$Arguments+='--whole-far'} + +for($Attempt=0;$Attempt -lt 2;$Attempt++) +{ + if($Attempt -gt 0) + { + Write-Host '0xC0000142 encountered; waiting 60 seconds before the single permitted retry.' + Start-Sleep -Seconds 60 + } + Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue + $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot ` + -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath + $PeakWorkingSet=0L;$PeakPrivateMemory=0L + while(-not$Child.HasExited) + { + $Child.Refresh() + if($Child.HasExited){break} + $Child.PriorityClass='BelowNormal' + $Child.ProcessorAffinity=[IntPtr]0xFF00 + $PeakWorkingSet=[Math]::Max($PeakWorkingSet,$Child.WorkingSet64) + $PeakPrivateMemory=[Math]::Max($PeakPrivateMemory,$Child.PrivateMemorySize64) + Start-Sleep -Milliseconds 100 + } + $Child.WaitForExit();$Child.Refresh() + if($Child.ExitCode -eq -1073741502 -and $Attempt -eq 0){continue} + if($Child.ExitCode -eq -1073741502){throw '0xC0000142 repeated after its single retry; stopping.'} + Write-Output "RUN_EXIT pid=$($Child.Id) priority=BelowNormal affinity=0xFF00 peak_working_set_gb=$([math]::Round($PeakWorkingSet/1GB,2)) peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) exit_code=$($Child.ExitCode) stdout=$StdoutPath stderr=$StderrPath" + exit $Child.ExitCode +} diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarViewer.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarViewer.ps1 new file mode 100644 index 0000000..c94ce61 --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/RunGpuFarViewer.ps1 @@ -0,0 +1,71 @@ +[CmdletBinding()] +param( + [ValidateSet('gpu','cpu','audit')][string]$Mode='gpu', + [string]$LogStem='gpu_far_viewer', + [int]$ClipRadius=3, + [string]$CaptureName='' +) + +$ErrorActionPreference='Stop' +if([IO.Path]::GetFileName($LogStem)-ne$LogStem){throw 'LogStem must be a filename stem.'} +$HostProcess=Get-Process -Id $PID +$HostProcess.PriorityClass='BelowNormal' +$HostProcess.ProcessorAffinity=[IntPtr]0xFF00 +$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..')) +$ExecutablePath=Join-Path $StandaloneRoot 'bin\VoxelForgeGpuFarHybridViewer.exe' +$StdoutPath=Join-Path $PSScriptRoot ($LogStem+'.stdout.log') +$StderrPath=Join-Path $PSScriptRoot ($LogStem+'.stderr.log') +if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "GPU far viewer not found: $ExecutablePath"} +$Arguments=@("--clip-radius=$ClipRadius",'--workers=8','--quit-after-clipmap') +if($Mode-eq'gpu'){$Arguments+='--gpu-far'} +elseif($Mode-eq'cpu'){$Arguments+='--cpu-far'} +else{$Arguments+='--gpu-far-audit'} +if($CaptureName) +{ + if([IO.Path]::GetFileName($CaptureName)-ne$CaptureName-or[IO.Path]::GetExtension($CaptureName)-ne'.png') + {throw 'CaptureName must be a PNG filename.'} + $Arguments+='--capture-after-clipmap' + $Arguments+=("--capture=Research/OwnerShapeWorld/gpu_field_port/$CaptureName") +} +$Command="`"$ExecutablePath`" $($Arguments -join ' ')" +Write-Output "VIEWER_COMMAND mode=$Mode $Command" + +for($Attempt=0;$Attempt -lt 2;$Attempt++) +{ + if($Attempt-gt 0) + { + Write-Host '0xC0000142 encountered; waiting 60 seconds before the single permitted retry.' + Start-Sleep -Seconds 60 + } + Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue + $Child=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot ` + -WindowStyle Hidden -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath + $PeakWorkingSet=0L;$PeakPrivateMemory=0L;$OverMemoryLimit=$false + while(-not$Child.HasExited) + { + $Child.Refresh() + if($Child.HasExited){break} + $Child.PriorityClass='BelowNormal' + $Child.ProcessorAffinity=[IntPtr]0xFF00 + $PeakWorkingSet=[Math]::Max($PeakWorkingSet,$Child.WorkingSet64) + $PeakPrivateMemory=[Math]::Max($PeakPrivateMemory,$Child.PrivateMemorySize64) + if($Child.PrivateMemorySize64-ge8GB) + { + $OverMemoryLimit=$true + $Child.Kill();$Child.WaitForExit();break + } + Start-Sleep -Milliseconds 100 + } + $Child.WaitForExit();$Child.Refresh() + if($Child.ExitCode-eq-1073741502-and$Attempt-eq0){continue} + if($Child.ExitCode-eq-1073741502){throw '0xC0000142 repeated after its single retry; stopping.'} + if($OverMemoryLimit){throw "Viewer stopped at the 8 GB per-process memory limit; stdout=$StdoutPath stderr=$StderrPath"} + $LogText=Get-Content -LiteralPath $StdoutPath -Raw + $Ready=$LogText.Contains('[InitialClipmapReady]') + $AuditComplete=($Mode-ne'audit')-or$LogText.Contains('[GpuFarSurvivalAudit]') + $AuditState=if($Mode-eq'audit'){if($AuditComplete){'reported'}else{'missing'}}else{'not-requested'} + Write-Output "RUN_EXIT pid=$($Child.Id) priority=BelowNormal affinity=0xFF00 peak_working_set_gb=$([math]::Round($PeakWorkingSet/1GB,2)) peak_private_gb=$([math]::Round($PeakPrivateMemory/1GB,2)) exit_code=$($Child.ExitCode) clipmap_ready=$Ready audit=$AuditState stdout=$StdoutPath stderr=$StderrPath" + if(-not$Ready){exit 4} + if(-not$AuditComplete){exit 5} + exit $Child.ExitCode +} diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/StartGpuFarViewer.ps1 b/Standalone/Research/OwnerShapeWorld/gpu_field_port/StartGpuFarViewer.ps1 new file mode 100644 index 0000000..4da3e0e --- /dev/null +++ b/Standalone/Research/OwnerShapeWorld/gpu_field_port/StartGpuFarViewer.ps1 @@ -0,0 +1,16 @@ +[CmdletBinding()] +param([int]$ClipRadius=3) + +$ErrorActionPreference='Stop' +$StandaloneRoot=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '..\..\..')) +$ExecutablePath=Join-Path $StandaloneRoot 'bin\VoxelForgeGpuFarHybridViewer.exe' +$StdoutPath=Join-Path $PSScriptRoot 'owner_hybrid_viewer.stdout.log' +$StderrPath=Join-Path $PSScriptRoot 'owner_hybrid_viewer.stderr.log' +if(-not(Test-Path -LiteralPath $ExecutablePath -PathType Leaf)){throw "GPU far viewer not found: $ExecutablePath"} +$Arguments=@('--gpu-far',"--clip-radius=$ClipRadius",'--workers=8') +Remove-Item -LiteralPath $StdoutPath,$StderrPath -Force -ErrorAction SilentlyContinue +$Viewer=Start-Process -FilePath $ExecutablePath -ArgumentList $Arguments -WorkingDirectory $StandaloneRoot ` + -WindowStyle Normal -PassThru -RedirectStandardOutput $StdoutPath -RedirectStandardError $StderrPath +$Viewer.PriorityClass='BelowNormal' +$Viewer.ProcessorAffinity=[IntPtr]0xFF00 +Write-Output "STARTED pid=$($Viewer.Id) priority=BelowNormal affinity=0xFF00 command=`"$ExecutablePath`" $($Arguments -join ' ') stdout=$StdoutPath stderr=$StderrPath" diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step3_cpu_far.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step3_cpu_far.png new file mode 100644 index 0000000..5756064 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step3_cpu_far.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_field_port/step3_gpu_far.png b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step3_gpu_far.png new file mode 100644 index 0000000..5756064 Binary files /dev/null and b/Standalone/Research/OwnerShapeWorld/gpu_field_port/step3_gpu_far.png differ diff --git a/Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl b/Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl index d3ad2e3..2f252fa 100644 --- a/Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl +++ b/Standalone/Research/OwnerShapeWorld/gpu_first_port/VoxelForgeGpuMesher.hlsl @@ -130,9 +130,10 @@ struct FarFieldSampleRange { uint offset; uint count; }; StructuredBuffer FarShapes : register(t10); StructuredBuffer FarSampleShapeRanges : register(t11); StructuredBuffer FarSampleShapeIds : register(t12); -StructuredBuffer FarRockAirSigns : register(t13); -StructuredBuffer FarProtectedAir : register(t14); +StructuredBuffer FarProtectedAir : register(t13); +StructuredBuffer FarProtectedSolid : register(t14); RWStructuredBuffer FarFieldDensity : register(u7); +StructuredBuffer GpuSignDensity : register(t15); cbuffer FarFieldConstants : register(b1) { @@ -186,31 +187,22 @@ void EvaluateFarFieldCS(uint3 dispatchId : SV_DispatchThreadID) const float base = min(position.z - FarStrataLo, FarStrataHi - position.z); float rockDistance = VF_FieldSmoothMin(base, max(cave, overlay), FarStrataBlend); const uint protectedAir = (FarProtectedAir[index >> 5] >> (index & 31)) & 1; - if (protectedAir != 0) rockDistance = min(rockDistance, -0.0625f); + const uint protectedSolid = (FarProtectedSolid[index >> 5] >> (index & 31)) & 1; const float rawDensity = -rockDistance; - const uint air = (FarRockAirSigns[index >> 5] >> (index & 31)) & 1; const float magnitude = max(abs(rawDensity), 0.05f); - FarFieldDensity[index] = air != 0 ? magnitude : -magnitude; + float density = rockDistance >= 0.0f ? -magnitude : magnitude; + if (protectedAir != 0) density = magnitude; + else if (protectedSolid != 0) density = -magnitude; + FarFieldDensity[index] = density; } groupshared uint ScanValues[256]; -[numthreads(256,1,1)] -void ClassifyCS(uint3 dispatchId : SV_DispatchThreadID, - uint3 groupId : SV_GroupID, - uint3 groupThreadId : SV_GroupThreadID) +void ClassifyAndScan(uint index, uint3 groupId, uint3 groupThreadId, uint caseIndex) { - const uint index = dispatchId.x; uint vertexCount = 0; - uint caseIndex = 0; if (index < CellCount) { - const uint3 c = uint3(index % CellsPerAxis, - (index / CellsPerAxis) % CellsPerAxis, - index / (CellsPerAxis * CellsPerAxis)); - [unroll] - for (uint corner = 0; corner < 8; ++corner) - caseIndex |= LatticeSign(Signs, int3(c) + CornerOffset[corner]) << corner; [unroll] for (uint slot = 0; slot < 15; ++slot) { @@ -235,6 +227,49 @@ void ClassifyCS(uint3 dispatchId : SV_DispatchThreadID, GroupVertexCounts[groupId.x] = ScanValues[255]; } +[numthreads(256,1,1)] +void ClassifyCS(uint3 dispatchId : SV_DispatchThreadID, + uint3 groupId : SV_GroupID, + uint3 groupThreadId : SV_GroupThreadID) +{ + const uint index = dispatchId.x; + uint caseIndex = 0; + if (index < CellCount) + { + const uint3 c = uint3(index % CellsPerAxis, + (index / CellsPerAxis) % CellsPerAxis, + index / (CellsPerAxis * CellsPerAxis)); + [unroll] + for (uint corner = 0; corner < 8; ++corner) + caseIndex |= LatticeSign(Signs, int3(c) + CornerOffset[corner]) << corner; + } + ClassifyAndScan(index, groupId, groupThreadId, caseIndex); +} + +uint GpuLatticeSign(int3 p) +{ + return GpuSignDensity[SampleIndex(p)] >= 0.0f ? 1 : 0; +} + +[numthreads(256,1,1)] +void ClassifyFarCS(uint3 dispatchId : SV_DispatchThreadID, + uint3 groupId : SV_GroupID, + uint3 groupThreadId : SV_GroupThreadID) +{ + const uint index = dispatchId.x; + uint caseIndex = 0; + if (index < CellCount) + { + const uint3 c = uint3(index % CellsPerAxis, + (index / CellsPerAxis) % CellsPerAxis, + index / (CellsPerAxis * CellsPerAxis)); + [unroll] + for (uint corner = 0; corner < 8; ++corner) + caseIndex |= GpuLatticeSign(int3(c) + CornerOffset[corner]) << corner; + } + ClassifyAndScan(index, groupId, groupThreadId, caseIndex); +} + StructuredBuffer InputGroupVertexCounts : register(t8); RWStructuredBuffer GroupOffsets : register(u3); RWByteAddressBuffer DrawArgs : register(u4); diff --git a/Standalone/Source/StandaloneViewer.cpp b/Standalone/Source/StandaloneViewer.cpp index feffab3..77292dc 100644 --- a/Standalone/Source/StandaloneViewer.cpp +++ b/Standalone/Source/StandaloneViewer.cpp @@ -2,6 +2,7 @@ #include "StreamingWorld.h" #include "StandaloneCoreScheduler.h" #include "../Core/src/StandaloneSnapshotAdapter.h" +#include "../Research/OwnerShapeWorld/gpu_field_port/GpuFarViewer.h" #include "VoxelDensityProfile.h" #ifdef TEXT @@ -23,6 +24,7 @@ #include #include #include +#include #include #include #include @@ -1345,6 +1347,16 @@ namespace std::chrono::steady_clock::time_point StartTime{}; bool bReported=false; }; + struct FGpuFarTile + { + ID3D11Buffer* VertexBuffer=nullptr; + ID3D11Buffer* IndexBuffer=nullptr; + ID3D11Buffer* IndirectArgs=nullptr; + ID3D11Query* CompletionEvent=nullptr; + uint64 VertexCapacityBytes=0; + double GpuDispatchMs=0.0; + bool bGpuComplete=false; + }; uint64 HashStandaloneFile(const std::string& Path) { @@ -1426,6 +1438,10 @@ namespace std::unordered_set DesiredTileSet; std::unordered_map,FStandaloneBatchKeyHash> BatchMembers; std::unordered_map StreamingBatches; + std::unordered_map GpuFarTiles; + std::deque GpuFarQueue; + std::unordered_set GpuFarQueued; + std::unordered_set GpuFarEligibleTiles; std::unordered_set DirtyBatchSet; std::unordered_map CrossingWork; std::vector DesiredTiles; @@ -1448,6 +1464,16 @@ namespace bool bShapeAuditOnly=false; bool bShapeCompare=false; bool bFarShapeStagePending=false; + bool bGpuFarEnabled=false; + bool bGpuFarAuditEnabled=false; + bool bGpuFarAuditReported=false; + std::unique_ptr GpuFarAudit; + bool bGpuFarFirstTileReported=false; + std::filesystem::path GpuFarShaderPath; + double GpuFarEnqueueWallSeconds=0.0; + uint64 GpuFarTilesSubmitted=0; + double GpuFarGpuDispatchMs=0.0; + double GpuFarGpuDispatchMaxMs=0.0; VoxelForgeCore::DetailVariant ShapeDetailVariant=VoxelForgeCore::DetailVariant::coherent; bool bWireframe=false; bool bCaptured=false; @@ -1513,6 +1539,17 @@ namespace if(Pair.second.IndexBuffer)Pair.second.IndexBuffer->Release(); } StreamingBatches.clear();BatchMembers.clear();ResidentTiles.clear(); + for(auto& Pair:GpuFarTiles) + { + if(Pair.second.VertexBuffer)Pair.second.VertexBuffer->Release(); + if(Pair.second.IndexBuffer)Pair.second.IndexBuffer->Release(); + if(Pair.second.IndirectArgs)Pair.second.IndirectArgs->Release(); + if(Pair.second.CompletionEvent)Pair.second.CompletionEvent->Release(); + } + GpuFarTiles.clear();GpuFarQueue.clear();GpuFarQueued.clear();GpuFarEligibleTiles.clear(); + GpuFarAudit.reset();bGpuFarAuditReported=false; + GpuFarTilesSubmitted=0;GpuFarEnqueueWallSeconds=0.0;GpuFarGpuDispatchMs=0.0; + GpuFarGpuDispatchMaxMs=0.0;bGpuFarFirstTileReported=false; DirtyBatchSet.clear();DesiredTiles.clear();DesiredTileSet.clear();CrossingWork.clear(); DeferredLeaverTiles.clear(); for(int32 I=0;I<=MaxLevels;++I) @@ -1668,7 +1705,7 @@ namespace cbuffer C:register(b0){float4x4 VP;float3 Camera;float FogStart;float3 Forward;float FogEnd;float3 LightDir;float Ambient;float4 Pad;} struct I{float3 p:POSITION;float3 n:NORMAL;}; struct O{float4 p:SV_POSITION;float3 world:TEXCOORD0;float3 n:NORMAL;}; -O VS(I i){O o;o.p=mul(float4(i.p,1),VP);o.world=i.p;o.n=normalize(i.n);return o;} + O VS(I i){O o;float3 p=i.p*Pad.x;o.p=mul(float4(p,1),VP);o.world=p;o.n=normalize(i.n);return o;} float4 PS(O i):SV_TARGET{ float3 N=normalize(i.n);float3 V=normalize(Camera-i.world);float3 L=normalize(-LightDir); float Diffuse=saturate(dot(N,L));float Head=saturate(dot(N,V)); @@ -1974,6 +2011,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; (!bShapeWorld&&MeshCache.IsLoaded())?&MeshCache:nullptr, bFarShapeStagePending?World->FullResClipLevels-1:-1, false,false,false,false,ShapeWorld.get()); + if(bGpuFarEnabled&&ShapeWorld&&ShapeWorld->hasFarWorld()) + { + ScheduleGpuFarTiles(DesiredTiles); + if(bGpuFarAuditEnabled) + { + std::vector AuditKeys;AuditKeys.reserve(GpuFarEligibleTiles.size()); + for(const FVoxelTileKey& Key:DesiredTiles) + if(GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key); + GpuFarAudit=std::make_unique(*World,*ShapeWorld,AuditKeys,false); + } + } if(bShapePipelineActive)Streamer->SetDesiredTileFilter(InitialNeighborhoodTiles); else if(CaptureTileLimit>0&&CapturePath.size()) Streamer->SetDesiredTileFilter(DesiredTiles); @@ -2025,6 +2073,17 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; Streamer.reset(); Streamer=std::make_unique(*World,WorkerCount,false,nullptr,-1, false,false,false,false,ShapeWorld.get()); + if(bGpuFarEnabled) + { + ScheduleGpuFarTiles(DesiredTiles); + if(bGpuFarAuditEnabled&&World) + { + std::vector AuditKeys;AuditKeys.reserve(GpuFarEligibleTiles.size()); + for(const FVoxelTileKey& Key:DesiredTiles) + if(GpuFarEligibleTiles.contains(Key))AuditKeys.push_back(Key); + GpuFarAudit=std::make_unique(*World,*ShapeWorld,AuditKeys,false); + } + } const FVector CameraVoxels(Position.X*4.0,Position.Y*4.0,Position.Z*4.0); ++CrossingIndex; const FStandaloneStreamingUpdate Update=Streamer->Start(CameraVoxels,ResidentKeys,CrossingIndex); @@ -2037,6 +2096,145 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; std::fflush(stdout); } + bool IsRenderedTileReady(const FVoxelTileKey& Key) + { + if(bGpuFarEnabled&&ShapeWorld&&ShapeWorld->hasFarWorld() + &&GpuFarEligibleTiles.contains(Key)) + { + auto It=GpuFarTiles.find(Key); + if(It==GpuFarTiles.end())return false; + FGpuFarTile& Tile=It->second; + if(Tile.bGpuComplete)return true; + if(Tile.CompletionEvent&&Context + &&Context->GetData(Tile.CompletionEvent,nullptr,0,D3D11_ASYNC_GETDATA_DONOTFLUSH)==S_OK) + { + Tile.bGpuComplete=true; + Tile.CompletionEvent->Release();Tile.CompletionEvent=nullptr; + } + return Tile.bGpuComplete; + } + return ResidentTiles.contains(Key); + } + + void ScheduleGpuFarTiles(const std::vector& Keys) + { + GpuFarQueue.clear();GpuFarQueued.clear();GpuFarEligibleTiles.clear(); + if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld())return; + std::vector CpuTiles;CpuTiles.reserve(Keys.size()); + for(const FVoxelTileKey& Key:Keys) + { + if(Key.LevelFullResClipLevels){CpuTiles.push_back(Key);continue;} + const VoxelForgeCore::TileKey CoreKey{{Key.Coord.X,Key.Coord.Y,Key.Coord.Z},Key.Level}; + if(ShapeWorld->isFarFieldTileEligible(CoreKey)) + { + GpuFarEligibleTiles.insert(Key); + if(!GpuFarTiles.contains(Key)) + {GpuFarQueue.push_back(Key);GpuFarQueued.insert(Key);} + } + else CpuTiles.push_back(Key); + } + if(Streamer)Streamer->SetDesiredTileFilter(CpuTiles); + const std::size_t CpuTransitionFallbacks=static_cast(std::count_if( + Keys.begin(),Keys.end(),[&](const FVoxelTileKey& Key) + {return Key.Level>=World->FullResClipLevels&&!GpuFarEligibleTiles.contains(Key);})); + std::printf("[GpuFarEligibility] eligible_gpu_tiles=%zu cpu_transition_fallback_tiles=%zu\n", + GpuFarEligibleTiles.size(),CpuTransitionFallbacks); + } + + void PumpGpuFarGeneration() + { + if(!bGpuFarEnabled||!ShapeWorld||!ShapeWorld->hasFarWorld()||!Device||!Context + ||GpuFarQueue.empty())return; + const auto Start=std::chrono::steady_clock::now(); + int32 Built=0; + do + { + const FVoxelTileKey Key=GpuFarQueue.front();GpuFarQueue.pop_front();GpuFarQueued.erase(Key); + if(!DesiredTileSet.contains(Key)||GpuFarTiles.contains(Key))continue; + FGpuFarTile Tile; + VoxelForgeGpuFarViewer::CpuMesh AuditMesh; + std::string Error;double EnqueueMs=0.0,GpuDispatchMs=0.0; + const VoxelForgeCore::TileKey CoreKey{{Key.Coord.X,Key.Coord.Y,Key.Coord.Z},Key.Level}; + if(!VoxelForgeGpuFarViewer::BuildFarTile(Device,Context,*ShapeWorld,CoreKey, + GpuFarShaderPath,&Tile.VertexBuffer,&Tile.IndexBuffer,&Tile.IndirectArgs, + &Tile.CompletionEvent, + (bGpuFarAuditEnabled&&Key.Level>=6)?&AuditMesh:nullptr, + &EnqueueMs,&GpuDispatchMs,Error)) + { + std::fprintf(stderr,"[GpuFarTileError] key=(%d,%d,%d,lod%d): %s\n", + Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,Error.c_str()); + bGpuFarEnabled=false; + break; + } + if(bGpuFarAuditEnabled&&GpuFarAudit&&Key.Level>=6) + { + FStandaloneCpuTile AuditTile;AuditTile.Key=Key; + AuditTile.SurfaceVertexCount=static_cast(AuditMesh.vertices.size()); + AuditTile.Vertices.reserve(AuditMesh.vertices.size()); + for(const auto& Vertex:AuditMesh.vertices) + AuditTile.Vertices.push_back({ + {Vertex.position[0]*0.01f,Vertex.position[1]*0.01f,Vertex.position[2]*0.01f}, + {Vertex.normal[0],Vertex.normal[1],Vertex.normal[2]}}); + AuditTile.Indices=std::move(AuditMesh.indices); + GpuFarAudit->AuditTile(AuditTile); + } + const uint64 Cells=static_cast(World->CoarseTileCells); + Tile.VertexCapacityBytes=Cells*Cells*Cells*15u*32u; + Tile.GpuDispatchMs=GpuDispatchMs; + GpuFarEnqueueWallSeconds+=EnqueueMs/1000.0; + GpuFarGpuDispatchMs+=GpuDispatchMs; + GpuFarGpuDispatchMaxMs=std::max(GpuFarGpuDispatchMaxMs,GpuDispatchMs); + ++GpuFarTilesSubmitted;++Built; + GpuFarTiles.emplace(Key,Tile); + if(!bGpuFarFirstTileReported) + { + bGpuFarFirstTileReported=true; + std::printf("[GpuFarFirstTile] key=(%d,%d,%d,lod%d) cpu_export_upload_enqueue_ms=%.3f vertex_capacity_mb=%.3f mode=direct-indirect-draw\n", + Key.Coord.X,Key.Coord.Y,Key.Coord.Z,Key.Level,EnqueueMs, + Tile.VertexCapacityBytes/(1024.0*1024.0)); + std::fflush(stdout); + } + } + while(Built<1&& !GpuFarQueue.empty() + &&std::chrono::duration(std::chrono::steady_clock::now()-Start).count()<1.0); + } + + void ToggleFarRenderer() + { + if(!bShapeWorld||!ShapeWorld||!ShapeWorld->hasFarWorld()) + { + bGpuFarEnabled=!bGpuFarEnabled; + std::printf("[FarRenderer] GPU far %s; it will take effect when far shapes are ready.\n", + bGpuFarEnabled?"enabled":"disabled"); + return; + } + bGpuFarEnabled=!bGpuFarEnabled; + if(bGpuFarEnabled) + { + ScheduleGpuFarTiles(DesiredTiles); + std::vector CpuFar; + for(const FVoxelTileKey& Key:DesiredTiles) + if(GpuFarEligibleTiles.contains(Key)&&ResidentTiles.contains(Key))CpuFar.push_back(Key); + for(const FVoxelTileKey& Key:CpuFar)RemoveResidentTile(Key); + } + else + { + GpuFarQueue.clear();GpuFarQueued.clear();GpuFarEligibleTiles.clear(); + Streamer->SetDesiredTileFilter(DesiredTiles); + } + std::vector ResidentKeys;ResidentKeys.reserve(ResidentTiles.size()); + for(const auto& Pair:ResidentTiles)ResidentKeys.push_back(Pair.first); + const FVector CameraVoxels(Position.X*4.0,Position.Y*4.0,Position.Z*4.0); + ++CrossingIndex; + const FStandaloneStreamingUpdate Update=Streamer->UpdateDesired( + CameraVoxels,ResidentKeys,CrossingIndex); + FCrossingWork& Work=CrossingWork[CrossingIndex]; + Work.SubmittedJobs=Update.NewJobs;Work.PendingJobs=Update.NewJobs; + Work.StartTime=std::chrono::steady_clock::now(); + std::printf("[FarRenderer] mode=%s cpu_far_jobs=%d gpu_far_ready=%zu gpu_far_queued=%zu (H toggles)\n", + bGpuFarEnabled?"GPU":"CPU",Update.NewJobs,GpuFarTiles.size(),GpuFarQueue.size()); + } + void TryCompleteShapePipeline() { if(!bShapePipelineActive||InitialNeighborhoodTiles.empty())return; @@ -2222,6 +2420,16 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; } DesiredTiles=NewDesired.Tiles; DesiredTileSet=std::move(NewSet); + for(auto It=GpuFarTiles.begin();It!=GpuFarTiles.end();) + { + if(DesiredTileSet.contains(It->first)){++It;continue;} + if(It->second.VertexBuffer)It->second.VertexBuffer->Release(); + if(It->second.IndexBuffer)It->second.IndexBuffer->Release(); + if(It->second.IndirectArgs)It->second.IndirectArgs->Release(); + if(It->second.CompletionEvent)It->second.CompletionEvent->Release(); + It=GpuFarTiles.erase(It); + } + ScheduleGpuFarTiles(DesiredTiles); RefreshResidentSkirtMasks(); for(int32 I=0;I<=MaxLevels;++I)DesiredByLevel[I]=0; for(const FVoxelTileKey& Tile:DesiredTiles)++DesiredByLevel[Tile.Level]; @@ -2489,7 +2697,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; { bool bDesiredResident=true; for(const FVoxelTileKey& Key:DesiredTiles) - if(!ResidentTiles.contains(Key)){bDesiredResident=false;break;} + if(!IsRenderedTileReady(Key)){bDesiredResident=false;break;} if(bDesiredResident) { std::vector ReadyToDrop; @@ -2531,7 +2739,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; const uint64 CacheHits=Streamer?Streamer->GetCacheHits():0; std::size_t DesiredResidentCount=0; for(const FVoxelTileKey& Key:DesiredTiles) - if(ResidentTiles.contains(Key))++DesiredResidentCount; + if(IsRenderedTileReady(Key))++DesiredResidentCount; std::sort(NearestReadyMilliseconds.begin(),NearestReadyMilliseconds.end()); const auto Percentile=[&](std::size_t Percent) { @@ -2568,7 +2776,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; } size_t DesiredResidentCount=0; for(const FVoxelTileKey& Key:DesiredTiles) - if(ResidentTiles.contains(Key))++DesiredResidentCount; + if(IsRenderedTileReady(Key))++DesiredResidentCount; if(!bInitialClipmapReported&&DesiredResidentCount==DesiredTiles.size()&&DirtyBatchSet.empty()) { bInitialClipmapReported=true; @@ -2576,10 +2784,28 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; std::printf("[InitialClipmapReady] %.3f s after cold process start; %zu tiles, %llu vertices, %llu indices, %zu GPU batches.\n", Elapsed,DesiredResidentCount,static_cast(TotalVertices), static_cast(TotalIndices),StreamingBatches.size()); + if(bGpuFarEnabled) + { + const double AverageGpuMs=GpuFarTilesSubmitted + ?GpuFarGpuDispatchMs/static_cast(GpuFarTilesSubmitted):0.0; + std::printf("[GpuFarTiming] submitted_tiles=%llu cpu_export_upload_enqueue_wall_ms=%.3f gpu_timestamp_total_ms=%.3f gpu_timestamp_mean_per_one_tile_frame_ms=%.6f gpu_timestamp_max_tile_ms=%.6f frame_budget_90hz_ms=11.111\n", + static_cast(GpuFarTilesSubmitted), + GpuFarEnqueueWallSeconds*1000.0,GpuFarGpuDispatchMs,AverageGpuMs,GpuFarGpuDispatchMaxMs); + } ReportTileTimings();std::fflush(stdout); if(bQuitAfterClipmap&&CapturePath.empty())PostMessageA(Window,WM_CLOSE,0,0); } - bGenerationRunning=DesiredResidentCount!=DesiredTiles.size()||!DirtyBatchSet.empty(); + if(bInitialClipmapReported&&bGpuFarAuditEnabled&&GpuFarAudit&&!bGpuFarAuditReported) + { + GpuFarAudit->Finalize();GpuFarAudit->Print(); + bGpuFarAuditReported=true; + std::printf("[GpuFarSurvivalAudit] mesh_source=GPU field_readback=required status=%s\n", + GpuFarAudit->Passing()?"PASS":"FAIL"); + std::fflush(stdout); + if(bQuitAfterClipmap&&CapturePath.empty())PostMessageA(Window,WM_CLOSE,0,0); + } + bGenerationRunning=DesiredResidentCount!=DesiredTiles.size()||!DirtyBatchSet.empty() + ||(bGpuFarEnabled&&!GpuFarQueue.empty()); } void CaptureRenderedFrame() @@ -2608,7 +2834,8 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; uint64 VisibleIndices=0;for(const auto& Pair:StreamingBatches)if(bShowLevel[Pair.first.Level])VisibleIndices+=Pair.second.IndexCount; std::ostringstream S;S<<"VoxelForge standalone | seed "<0&&bFullDetailReported&&!bFlightStarted) { @@ -2658,6 +2886,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; C.ViewProjection=Transpose(Combined);C.Camera[0]=Position.X;C.Camera[1]=Position.Y;C.Camera[2]=Position.Z; C.Forward[0]=F.X;C.Forward[1]=F.Y;C.Forward[2]=F.Z;C.LightDirection[0]=-0.5f;C.LightDirection[1]=-0.72f;C.LightDirection[2]=-0.48f; C.FogStart=260.0f;C.FogEnd=2600.0f;C.Ambient=0.36f; + C.Padding[0]=1.0f; Context->UpdateSubresource(Constants,0,nullptr,&C,0,0);Context->VSSetConstantBuffers(0,1,&Constants);Context->PSSetConstantBuffers(0,1,&Constants); Context->VSSetShader(VertexShader,nullptr,0);Context->PSSetShader(PixelShader,nullptr,0); for(const auto& Pair:StreamingBatches) @@ -2667,6 +2896,21 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; const UINT Stride=sizeof(FGpuVertex),Offset=0;Context->IASetVertexBuffers(0,1,&L.VertexBuffer,&Stride,&Offset); Context->IASetIndexBuffer(L.IndexBuffer,DXGI_FORMAT_R32_UINT,0);Context->DrawIndexed(L.IndexCount,0,0); } + if(bGpuFarEnabled&&ShapeWorld&&ShapeWorld->hasFarWorld()) + { + C.Padding[0]=0.01f; + Context->UpdateSubresource(Constants,0,nullptr,&C,0,0); + Context->VSSetConstantBuffers(0,1,&Constants);Context->PSSetConstantBuffers(0,1,&Constants); + for(const auto& Pair:GpuFarTiles) + { + const FGpuFarTile& Tile=Pair.second; + if(!bShowLevel[Pair.first.Level]||!Tile.VertexBuffer||!Tile.IndexBuffer||!Tile.IndirectArgs)continue; + const UINT Stride=32,Offset=0; + Context->IASetVertexBuffers(0,1,&Tile.VertexBuffer,&Stride,&Offset); + Context->IASetIndexBuffer(Tile.IndexBuffer,DXGI_FORMAT_R32_UINT,0); + Context->DrawIndexedInstancedIndirect(Tile.IndirectArgs,0); + } + } CaptureRenderedFrame(); SwapChain->Present(1,0);UpdateTitle(); } @@ -2693,6 +2937,7 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; std::printf("Round-tunnel comparison: %s wall field.\n", GApp->bRoundTunnelEnhanced?"detailed/smoothed":"legacy smooth tube"); } + if(!bWasDown&&W=='H')GApp->ToggleFarRenderer(); if(!bWasDown&&(W=='N'||W=='P'))GApp->RequestSeed(W=='N'?1:-1); if(!bWasDown&&W>='1'&&W<='8'){const int32 I=static_cast(W-'1');GApp->bShowLevel[I]=!GApp->bShowLevel[I];} if(!bWasDown&&W=='9')GApp->bShowLevel[8]=!GApp->bShowLevel[8]; @@ -4425,8 +4670,15 @@ float3 Lit=Rock*(Ambient+0.88*Diffuse)+float3(0.50,0.69,0.82)*Head*0.75; int main(int Argc,char** Argv) { + SetPriorityClass(GetCurrentProcess(),BELOW_NORMAL_PRIORITY_CLASS); + SetProcessAffinityMask(GetCurrentProcess(),static_cast(0xFF00)); const auto ColdProcessStart=std::chrono::steady_clock::now(); FApp App;App.DataPath=DefaultDataPath().string();App.ColdProcessStart=ColdProcessStart; + wchar_t ExecutablePath[MAX_PATH]{}; + const DWORD ExecutablePathLength=GetModuleFileNameW(nullptr,ExecutablePath,MAX_PATH); + if(ExecutablePathLength>0&&ExecutablePathLength