streaming: split dense passage-junction grid fills; heavy-first order + far-parent deferral in Unreal (L85)

Tiles with passage-overlap score >= 9 split grid fill over up to 4 helper lanes (disjoint rows,
same per-sample order, join before MC) when workers are idle. FI junction tile fill 173 -> 99 ms;
estimated serial floor ~0.55-0.60 s. Readiness essentially unchanged (FI/TN LOD0-2 ~2.15 s @W12):
the limit is now aggregate work, not single tiles.
Unreal: LOD0-2 ordered by passage overlap before travel/distance (L80 modes kept); far parents
built after LOD0-1, LOD2 gated on them; helper budget via ParallelFor. Unreal null-RHI FI
arrival: LOD0-2 resident at 3.65 s.
All tile hashes identical to L84 (four whole views). Fast tier PASS; Linux UE host builds.

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

Standalone/Source/VoxelForge/Private/VoxelMarchingCubesMesher.cpp

@@ -17,6 +17,9 @@
 #include "HAL/IConsoleManager.h"
 #include "Misc/CommandLine.h"
 #include <atomic>
+#include <array>
+#include <algorithm>
+#include <iterator>
 #include <bit>
 #include <cstdlib>
 #include <immintrin.h>
@@ -24,6 +27,69 @@
 
 namespace
 {
+    struct FVoxelMesherThreadState
+    {
+        int32 OctaveBias=0,SampleStep=1,TileCells=0,TileReachBlocks=0,TileReachBlockIndex=INDEX_NONE;
+        FIntVector TileOrigin=FIntVector::ZeroValue;
+        uint8 TilePostReachFlags=VoxelGenLOD::AllTileReach;
+        bool bTilePostReachDebug=false,bTilePostReachBypass=false;
+        bool bTileReachDiagnosticTile=false,bTileCoreReachProofEnabled=false;
+        bool bTileCoreReachDecisionRecorded=false,bTileBlockReachValid=false;
+        float TileReachScale=1.0f;
+        std::array<uint8,VoxelGenLOD::TileReachMaxBlockCount> TileBlockReachFlags{};
+
+        static FVoxelMesherThreadState Capture()
+        {
+            FVoxelMesherThreadState State;
+            State.OctaveBias=VoxelGenLOD::OctaveBias;
+            State.SampleStep=VoxelGenLOD::SampleStep;
+            State.TileOrigin=VoxelGenLOD::TileOriginVoxels;
+            State.TileCells=VoxelGenLOD::TileCellsPerAxis;
+            State.TilePostReachFlags=VoxelGenLOD::TilePostReachFlags;
+            State.bTilePostReachDebug=VoxelGenLOD::bTilePostReachDebug;
+            State.bTilePostReachBypass=VoxelGenLOD::bTilePostReachBypass;
+            State.bTileReachDiagnosticTile=VoxelGenLOD::bTileReachDiagnosticTile;
+            State.bTileCoreReachProofEnabled=VoxelGenLOD::bTileCoreReachProofEnabled;
+            State.bTileCoreReachDecisionRecorded=VoxelGenLOD::bTileCoreReachDecisionRecorded;
+            State.bTileBlockReachValid=VoxelGenLOD::bTileBlockReachValid;
+            State.TileReachScale=VoxelGenLOD::TileReachScale;
+            State.TileReachBlocks=VoxelGenLOD::TileReachBlocksPerAxis;
+            State.TileReachBlockIndex=VoxelGenLOD::TileReachBlockIndex;
+            std::copy(std::begin(VoxelGenLOD::TileBlockReachFlags),
+                std::end(VoxelGenLOD::TileBlockReachFlags),State.TileBlockReachFlags.begin());
+            return State;
+        }
+
+        void Apply() const
+        {
+            VoxelGenLOD::OctaveBias=OctaveBias;
+            VoxelGenLOD::SampleStep=SampleStep;
+            VoxelGenLOD::TileOriginVoxels=TileOrigin;
+            VoxelGenLOD::TileCellsPerAxis=TileCells;
+            VoxelGenLOD::TilePostReachFlags=TilePostReachFlags;
+            VoxelGenLOD::bTilePostReachDebug=bTilePostReachDebug;
+            VoxelGenLOD::bTilePostReachBypass=bTilePostReachBypass;
+            VoxelGenLOD::bTileReachDiagnosticTile=bTileReachDiagnosticTile;
+            VoxelGenLOD::bTileCoreReachProofEnabled=bTileCoreReachProofEnabled;
+            VoxelGenLOD::bTileCoreReachDecisionRecorded=bTileCoreReachDecisionRecorded;
+            VoxelGenLOD::bTileBlockReachValid=bTileBlockReachValid;
+            VoxelGenLOD::TileReachScale=TileReachScale;
+            VoxelGenLOD::TileReachBlocksPerAxis=TileReachBlocks;
+            VoxelGenLOD::TileReachBlockIndex=TileReachBlockIndex;
+            std::copy(TileBlockReachFlags.begin(),TileBlockReachFlags.end(),
+                std::begin(VoxelGenLOD::TileBlockReachFlags));
+        }
+    };
+
+    struct FVoxelMesherThreadStateScope
+    {
+        explicit FVoxelMesherThreadStateScope(const FVoxelMesherThreadState& InTarget)
+            :Previous(FVoxelMesherThreadState::Capture())
+        { InTarget.Apply(); }
+        ~FVoxelMesherThreadStateScope(){Previous.Apply();}
+        FVoxelMesherThreadState Previous;
+    };
+
     int32 GVoxelForgeUseOperatorBlock = 1;
     FAutoConsoleVariableRef CVarVoxelForgeUseOperatorBlock(
         TEXT("voxel.UseOperatorBlock"),
@@ -957,6 +1023,7 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
     // capacité, donc plus de malloc/free de ~170 Ko (35³ floats) par tuile.
     static thread_local TArray<float> DensityGrid;
     DensityGrid.SetNumUninitialized(MDim * MDim * MDim);
+    float* const DensityGridData=DensityGrid.GetData();
     const FVoxelSharedDensityGrid* ReuseGrid = InSharedDensityGrid != nullptr
         ? InSharedDensityGrid : SharedDensityGrid;
     const int32 SharedDeltaX = ReuseGrid
@@ -1063,14 +1130,7 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
     const bool bUseOperatorBlock = !bUseSharedDensityGrid
         && !bUseParameterizedBrickEvaluator
         && GVoxelForgeUseOperatorBlock != 0;
-    if (bUseOperatorBlock)
-    {
-        // The grid includes one sample of halo on every side.  The generator groups these exact
-        // lattice points by chunk key and lets the prepared op stack fill each group once.
-        Generator->BeginDensityBlock(
-            OriginVoxels + FIntVector(-Step, -Step, GzLo * Step),
-            Step, MDim, MDim, GzHi - GzLo + 1);
-    }
+    bool bGridFillParallel=false;
     if (bUseParameterizedBrickEvaluator)
     {
 #if WITH_EDITOR
@@ -1238,7 +1298,9 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
             * (VoxelGenLOD::IsUseFusedEvaluatorEnabled() ? 4 : 1);
         const auto XGroups = BuildAxisGroups(OriginVoxels.X,RoomCacheRegionSize);
         const auto YGroups = BuildAxisGroups(OriginVoxels.Y,RoomCacheRegionSize);
-        auto SampleGridPoint = [&](int32 gx, int32 gy, int32 gz)
+        auto SampleGridPoint = [&](int32 gx, int32 gy, int32 gz,
+                                  FVoxelMarchingCubesStageTimings* StageTimings,
+                                  int64* DensitySampleCount,float* GridData)
         {
             const int32 LocalIndex = ((gz + 1) * MDim + (gy + 1)) * MDim + (gx + 1);
             if(bSkipProofBlockGrid&&RequiredGridSamples[LocalIndex]==0)return;
@@ -1247,7 +1309,7 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
                 const int32 SharedX = SharedDeltaX + gx;
                 const int32 SharedY = SharedDeltaY + gy;
                 const int32 SharedZ = SharedDeltaZ + gz;
-                DensityGrid[LocalIndex] = ReuseGrid->Samples[
+                GridData[LocalIndex] = ReuseGrid->Samples[
                     (SharedZ * ReuseGrid->Dim + SharedY) * ReuseGrid->Dim + SharedX];
             }
             else
@@ -1272,22 +1334,22 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
                     bBoundarySample
                         ?VoxelForgeCore::floating_island_detail::EIslandFieldCaller::Boundary
                         :VoxelForgeCore::floating_island_detail::EIslandFieldCaller::GridFill);
-                DensityGrid[LocalIndex] = ShapeWorldContext
+                GridData[LocalIndex] = ShapeWorldContext
                     ?ShapeWorldContext->DensityAtLattice(SampleX,SampleY,SampleZ)
                     :Generator->GetDensityAtLattice(SampleX,SampleY,SampleZ,
                         bCapturePreEditGrid ? &PreEditDensity : nullptr);
-                if(OutStageTimings)
+                if(StageTimings)
                 {
-                    if(bBoundarySample)++OutStageTimings->DensityGridHaloSamples;
-                    else ++OutStageTimings->DensityGridInteriorSamples;
+                    if(bBoundarySample)++StageTimings->DensityGridHaloSamples;
+                    else ++StageTimings->DensityGridInteriorSamples;
                 }
                 if (bCapturePreEditGrid)
                 {
                     OutPreEditDensityGrid->Samples[LocalIndex] = PreEditDensity;
                 }
-                if (OutDensitySampleCount)
+                if (DensitySampleCount)
                 {
-                    ++(*OutDensitySampleCount);
+                    ++(*DensitySampleCount);
                 }
             }
         };
@@ -1308,7 +1370,9 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
         const bool bUseShapeWorldXBatch=ActiveShapeWorld!=nullptr
             &&!bUseSharedDensityGrid&&!bCapturePreEditGrid
             &&(!bSkipProofBlockGrid||bStandaloneEnableActiveLaneBatching);
-        const auto SampleShapeWorldXBatch=[&](int32 GX0,int32 BatchCount,int32 gy,int32 gz)
+        const auto SampleShapeWorldXBatch=[&](int32 GX0,int32 BatchCount,int32 gy,int32 gz,
+            FVoxelMarchingCubesStageTimings* StageTimings,int64* DensitySampleCount,
+            float* GridData)
         {
             float Densities[8]{};
             const bool bBoundarySample=GX0<0||GX0+BatchCount-1>GridDim
@@ -1320,24 +1384,24 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
                     :VoxelForgeCore::floating_island_detail::EIslandFieldCaller::GridFill);
             ActiveShapeWorld->DensityAtLatticeX8(OriginVoxels.X+GX0*Step,Step,
                 OriginVoxels.Y+gy*Step,OriginVoxels.Z+gz*Step,BatchCount,Densities);
-            if(OutStageTimings)
+            if(StageTimings)
             {
-                ++OutStageTimings->X8DensityCalls;
-                OutStageTimings->X8ActiveLanes+=static_cast<uint64>(BatchCount);
+                ++StageTimings->X8DensityCalls;
+                StageTimings->X8ActiveLanes+=static_cast<uint64>(BatchCount);
                 for(int32 Lane=0;Lane<BatchCount;++Lane)
                 {
                     const int32 SampleGX=GX0+Lane;
                     if(SampleGX<0||SampleGX>GridDim||gy<0||gy>GridDim
                         ||gz<0||gz>GridDim)
-                        ++OutStageTimings->DensityGridHaloSamples;
-                    else ++OutStageTimings->DensityGridInteriorSamples;
+                        ++StageTimings->DensityGridHaloSamples;
+                    else ++StageTimings->DensityGridInteriorSamples;
                 }
             }
             for(int32 Lane=0;Lane<BatchCount;++Lane)
             {
                 const int32 GX=GX0+Lane;
                 const int32 LocalIndex=((gz+1)*MDim+(gy+1))*MDim+(GX+1);
-                DensityGrid[LocalIndex]=Densities[Lane];
+                GridData[LocalIndex]=Densities[Lane];
             }
             const int32 LastGX=GX0+BatchCount-1;
             const int32 ReachBlockX=FMath::Clamp(LastGX,0,CellsPerAxis-1)
@@ -1348,9 +1412,202 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
                 /VoxelGenLOD::TileReachBlockCells;
             VoxelGenLOD::TileReachBlockIndex=
                 ((ReachBlockZ*ReachBlocksPerAxis)+ReachBlockY)*ReachBlocksPerAxis+ReachBlockX;
-            if(OutDensitySampleCount)*OutDensitySampleCount+=BatchCount;
+            if(DensitySampleCount)*DensitySampleCount+=BatchCount;
         };
-        if(bUseShapeWorldZBatch&&!bUseVerticalCacheTraversal)
+        const int32 GridZSampleCount=GzHi-GzLo+1;
+        const int32 GridYSampleCount=GridDim+2;
+        const bool bParallelizeGridRows=bUseShapeWorldZBatch
+            ||bUseVerticalCacheTraversal||!bUseShapeWorldXBatch;
+        const std::size_t GridFillLaneCount=std::min<std::size_t>({
+            GridFillExecution.WorkerLimit,16,
+            static_cast<std::size_t>(FMath::Max(0,
+                bParallelizeGridRows?GridYSampleCount:GridZSampleCount))});
+        const bool bUseParallelShapeWorldGrid=GridFillExecution.ParallelFor!=nullptr
+            &&GridFillLaneCount>1&&GridFillExecution.BeginWorkerContext!=nullptr
+            &&GridFillExecution.EndWorkerContext!=nullptr
+            &&!bUseSharedDensityGrid&&!bCapturePreEditGrid&&!bSkipProofBlockGrid
+            &&OutCaptureGrid==nullptr;
+        if(bUseOperatorBlock&&!bUseParallelShapeWorldGrid)
+        {
+            // The grid includes one sample of halo on every side. The prepared operator block
+            // is thread-local, so a parallel fill below opens one block per worker slab.
+            Generator->BeginDensityBlock(
+                OriginVoxels+FIntVector(-Step,-Step,GzLo*Step),
+                Step,MDim,MDim,GzHi-GzLo+1);
+        }
+        if(bUseParallelShapeWorldGrid)
+        {
+            bGridFillParallel=true;
+            const FVoxelMesherThreadState TileThreadState=FVoxelMesherThreadState::Capture();
+            std::array<int64,16> LaneSampleCounts{};
+            std::array<FVoxelMarchingCubesStageTimings,16> LaneTimings{};
+            std::atomic<bool> bAborted{false};
+            auto RunLane=[&](std::size_t Lane,std::size_t Worker)
+            {
+                void* WorkerState=GridFillExecution.BeginWorkerContext(
+                    GridFillExecution.WorkerContext,Worker);
+                int64& LaneSamples=LaneSampleCounts[Lane];
+                FVoxelMarchingCubesStageTimings& StageTimings=LaneTimings[Lane];
+                {
+                    FVoxelMesherThreadStateScope ThreadStateScope(TileThreadState);
+                    if(bUseOperatorBlock)
+                        Generator->BeginDensityBlock(
+                            OriginVoxels+FIntVector(-Step,-Step,GzLo*Step),Step,
+                            MDim,MDim,GridZSampleCount);
+                    if(bUseShapeWorldZBatch)
+                    {
+                        const int32 FirstGy=-1+static_cast<int32>(
+                            (static_cast<std::size_t>(GridYSampleCount)*Lane)/GridFillLaneCount);
+                        const int32 EndGy=-1+static_cast<int32>(
+                            (static_cast<std::size_t>(GridYSampleCount)*(Lane+1))/GridFillLaneCount);
+                        for(int32 gy=FirstGy;gy<EndGy;++gy)
+                        {
+                            if(ShouldAbortWork())
+                            {bAborted.store(true,std::memory_order_relaxed);break;}
+                            for(const std::vector<int32>& XGroup:XGroups)
+                            for(const int32 gx:XGroup)
+                            {
+                                if(ShouldAbortWork())
+                                {
+                                    bAborted.store(true,std::memory_order_relaxed);
+                                    break;
+                                }
+                                float Densities[CHUNK_SIZE+3]{};
+                                const bool bBoundarySample=gx<0||gx>GridDim
+                                    ||gy<0||gy>GridDim||GzLo<0||GzHi>GridDim;
+                                VoxelForgeCore::Linux77Profiler::FScopedHalo l77Halo(
+                                    bBoundarySample);
+                                VoxelForgeCore::floating_island_detail::FScopedIslandFieldCaller
+                                    islandCaller(bBoundarySample
+                                        ?VoxelForgeCore::floating_island_detail::EIslandFieldCaller::Boundary
+                                        :VoxelForgeCore::floating_island_detail::EIslandFieldCaller::GridFill);
+                                ActiveShapeWorld->DensityAtLatticeZRun(
+                                    OriginVoxels.X+gx*Step,OriginVoxels.Y+gy*Step,
+                                    OriginVoxels.Z+GzLo*Step,Step,GridZSampleCount,Densities);
+                                for(int32 ZLane=0;ZLane<GridZSampleCount;++ZLane)
+                                {
+                                    const int32 gz=GzLo+ZLane;
+                                    const int32 LocalIndex=
+                                        ((gz+1)*MDim+(gy+1))*MDim+(gx+1);
+                                    DensityGridData[LocalIndex]=Densities[ZLane];
+                                    if(gx<0||gx>GridDim||gy<0||gy>GridDim
+                                        ||gz<0||gz>GridDim)
+                                        ++StageTimings.DensityGridHaloSamples;
+                                    else ++StageTimings.DensityGridInteriorSamples;
+                                }
+                                LaneSamples+=GridZSampleCount;
+                                const int32 ReachBlockX=FMath::Clamp(gx,0,CellsPerAxis-1)
+                                    /VoxelGenLOD::TileReachBlockCells;
+                                const int32 ReachBlockY=FMath::Clamp(gy,0,CellsPerAxis-1)
+                                    /VoxelGenLOD::TileReachBlockCells;
+                                const int32 ReachBlockZ=FMath::Clamp(GzHi,0,CellsPerAxis-1)
+                                    /VoxelGenLOD::TileReachBlockCells;
+                                VoxelGenLOD::TileReachBlockIndex=
+                                    ((ReachBlockZ*ReachBlocksPerAxis)+ReachBlockY)
+                                        *ReachBlocksPerAxis+ReachBlockX;
+                            }
+                            if(bAborted.load(std::memory_order_relaxed))break;
+                        }
+                    }
+                    else if(bUseVerticalCacheTraversal||!bUseShapeWorldXBatch)
+                    {
+                        const int32 FirstGy=-1+static_cast<int32>(
+                            (static_cast<std::size_t>(GridYSampleCount)*Lane)/GridFillLaneCount);
+                        const int32 EndGy=-1+static_cast<int32>(
+                            (static_cast<std::size_t>(GridYSampleCount)*(Lane+1))/GridFillLaneCount);
+                        for(int32 gy=FirstGy;gy<EndGy;++gy)
+                        {
+                            if(ShouldAbortWork())
+                            {bAborted.store(true,std::memory_order_relaxed);break;}
+                            for(const std::vector<int32>& XGroup:XGroups)
+                            for(const int32 gx:XGroup)
+                            {
+                                for(int32 gz=GzLo;gz<=GzHi;++gz)
+                                {
+                                    if(ShouldAbortWork())
+                                    {bAborted.store(true,std::memory_order_relaxed);break;}
+                                    SampleGridPoint(gx,gy,gz,&StageTimings,&LaneSamples,
+                                        DensityGridData);
+                                }
+                                if(bAborted.load(std::memory_order_relaxed))break;
+                            }
+                            if(bAborted.load(std::memory_order_relaxed))break;
+                        }
+                    }
+                    else
+                    {
+                        const int32 FirstGz=GzLo+static_cast<int32>(
+                            (static_cast<std::size_t>(GridZSampleCount)*Lane)/GridFillLaneCount);
+                        const int32 EndGz=GzLo+static_cast<int32>(
+                            (static_cast<std::size_t>(GridZSampleCount)*(Lane+1))/GridFillLaneCount);
+                        for(int32 gz=FirstGz;gz<EndGz;++gz)
+                        {
+                            if(ShouldAbortWork())
+                            {bAborted.store(true,std::memory_order_relaxed);break;}
+                            for(const std::vector<int32>& YGroup:YGroups)
+                            for(const std::vector<int32>& XGroup:XGroups)
+                            for(const int32 gy:YGroup)
+                            {
+                                if(ShouldAbortWork())
+                                {bAborted.store(true,std::memory_order_relaxed);break;}
+                                for(std::size_t Offset=0;Offset<XGroup.size();)
+                                {
+                                    const int32 BatchCount=static_cast<int32>(std::min<std::size_t>(
+                                        8,XGroup.size()-Offset));
+                                    const int32 GX0=XGroup[Offset];
+                                    bool bContiguous=true;
+                                    for(int32 BatchLane=1;BatchLane<BatchCount;++BatchLane)
+                                        bContiguous&=XGroup[
+                                            Offset+static_cast<std::size_t>(BatchLane)]==GX0+BatchLane;
+                                    if(!bContiguous)
+                                    {
+                                        for(int32 BatchLane=0;BatchLane<BatchCount;++BatchLane)
+                                            SampleGridPoint(XGroup[
+                                                Offset+static_cast<std::size_t>(BatchLane)],gy,gz,
+                                                &StageTimings,&LaneSamples,DensityGridData);
+                                    }
+                                    else
+                                        SampleShapeWorldXBatch(GX0,BatchCount,gy,gz,
+                                            &StageTimings,&LaneSamples,DensityGridData);
+                                    Offset+=static_cast<std::size_t>(BatchCount);
+                                }
+                            }
+                            if(bAborted.load(std::memory_order_relaxed))break;
+                        }
+                    }
+                    if(bUseOperatorBlock)Generator->EndDensityBlock();
+                }
+                GridFillExecution.EndWorkerContext(
+                    GridFillExecution.WorkerContext,WorkerState,Worker);
+            };
+            struct FDispatchContext{decltype(RunLane)* Run=nullptr;};
+            FDispatchContext DispatchContext{&RunLane};
+            const auto DispatchLane=[](void* RawContext,std::size_t Lane,std::size_t Worker)
+            {
+                auto& Context=*static_cast<FDispatchContext*>(RawContext);
+                (*Context.Run)(Lane,Worker);
+            };
+            GridFillExecution.ParallelFor(GridFillExecution.SchedulerContext,
+                GridFillLaneCount,GridFillLaneCount,&DispatchContext,DispatchLane);
+            if(OutDensitySampleCount)
+                for(std::size_t Lane=0;Lane<GridFillLaneCount;++Lane)
+                    *OutDensitySampleCount+=LaneSampleCounts[Lane];
+            if(OutStageTimings)
+            {
+                OutStageTimings->GridFillWorkerCount=static_cast<uint32>(GridFillLaneCount);
+                for(std::size_t Lane=0;Lane<GridFillLaneCount;++Lane)
+                {
+                    OutStageTimings->DensityGridInteriorSamples+=
+                        LaneTimings[Lane].DensityGridInteriorSamples;
+                    OutStageTimings->DensityGridHaloSamples+=
+                        LaneTimings[Lane].DensityGridHaloSamples;
+                    OutStageTimings->X8DensityCalls+=LaneTimings[Lane].X8DensityCalls;
+                    OutStageTimings->X8ActiveLanes+=LaneTimings[Lane].X8ActiveLanes;
+                }
+            }
+            if(bAborted.load(std::memory_order_relaxed))return FVoxelMeshData();
+        }
+        else if(bUseShapeWorldZBatch&&!bUseVerticalCacheTraversal)
         {
             for(const std::vector<int32>& YGroup:YGroups)
             for(const std::vector<int32>& XGroup:XGroups)
@@ -1419,7 +1676,8 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
                         if (bUseOperatorBlock) { Generator->EndDensityBlock(); }
                         return FVoxelMeshData();
                     }
-                    SampleGridPoint(gx,gy,gz);
+                    SampleGridPoint(gx,gy,gz,OutStageTimings,OutDensitySampleCount,
+                        DensityGridData);
                 }
             }
         }
@@ -1460,8 +1718,11 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
                                     if(RequiredGridSamples[NextIndex]==0)break;
                                     ++BatchCount;
                                 }
-                                if(BatchCount<2)SampleGridPoint(GX0,gy,gz);
-                                else SampleShapeWorldXBatch(GX0,BatchCount,gy,gz);
+                                if(BatchCount<2)SampleGridPoint(
+                                    GX0,gy,gz,OutStageTimings,OutDensitySampleCount,
+                                    DensityGridData);
+                                else SampleShapeWorldXBatch(GX0,BatchCount,gy,gz,
+                                    OutStageTimings,OutDensitySampleCount,DensityGridData);
                                 Offset+=static_cast<size_t>(BatchCount);
                                 continue;
                             }
@@ -1474,18 +1735,21 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
                             if(!bContiguous)
                             {
                                 for(int32 Lane=0;Lane<BatchCount;++Lane)
-                                    SampleGridPoint(XGroup[Offset+static_cast<size_t>(Lane)],gy,gz);
+                                    SampleGridPoint(XGroup[Offset+static_cast<size_t>(Lane)],gy,gz,
+                                        OutStageTimings,OutDensitySampleCount,DensityGridData);
                                 Offset+=static_cast<size_t>(BatchCount);
                                 continue;
                             }
 
-                            SampleShapeWorldXBatch(GX0,BatchCount,gy,gz);
+                            SampleShapeWorldXBatch(GX0,BatchCount,gy,gz,
+                                OutStageTimings,OutDensitySampleCount,DensityGridData);
                             Offset+=static_cast<size_t>(BatchCount);
                         }
                     }
                     else
                     {
-                        for(const int32 gx:XGroup)SampleGridPoint(gx,gy,gz);
+                        for(const int32 gx:XGroup)SampleGridPoint(
+                            gx,gy,gz,OutStageTimings,OutDensitySampleCount,DensityGridData);
                     }
                 }
             }
@@ -1506,7 +1770,7 @@ FVoxelMeshData UVoxelMarchingCubesMesher::GenerateMesh(FIntVector OriginVoxels,
                 GridSampleCount * sizeof(float));
         }
     }
-    if (bUseOperatorBlock)
+    if (bUseOperatorBlock&&!bGridFillParallel)
     {
         Generator->EndDensityBlock();
     }

Standalone/Source/VoxelForge/Public/VoxelMarchingCubesMesher.h

@@ -13,6 +13,7 @@
 
 #include "CoreMinimal.h"
 #include <atomic>
+#include <cstddef>
 #include <vector>
 #include "VoxelTypes.h"   // Pour FVoxelMeshData, CHUNK_SIZE, VOXEL_SIZE, etc.
 #include "VoxelGenerator.h"
@@ -41,10 +42,29 @@ struct FVoxelSharedDensityGrid
     }
 };
 
+// Host-owned scheduling hooks for a predicted-heavy density-grid fill. Each task owns one
+// deterministic Z slab and gets the source tile's worker-local sampling context.
+struct FVoxelGridFillExecution
+{
+    using FTask = void(*)(void* taskContext, std::size_t index, std::size_t workerIndex);
+    using FParallelFor = double(*)(void* schedulerContext, std::size_t count,
+        std::size_t workerLimit, void* taskContext, FTask task);
+    using FBeginWorkerContext = void*(*)(void* context, std::size_t workerIndex);
+    using FEndWorkerContext = void(*)(void* context, void* workerState,
+        std::size_t workerIndex);
+    void* SchedulerContext = nullptr;
+    FParallelFor ParallelFor = nullptr;
+    std::size_t WorkerLimit = 1;
+    void* WorkerContext = nullptr;
+    FBeginWorkerContext BeginWorkerContext = nullptr;
+    FEndWorkerContext EndWorkerContext = nullptr;
+};
+
 // Optional, sampled timing output for the research audit; null in production calls.
 // Mesures echantillonnees facultatives pour l'audit; pointeur nul en production.
 struct FVoxelMarchingCubesStageTimings
 {
+    uint32 GridFillWorkerCount = 1;
     double DensityGridMilliseconds = 0.0;
     double CellClassificationMilliseconds = 0.0;
     double GradientNormalsMilliseconds = 0.0;
@@ -161,6 +181,8 @@ public:
     // GenerateMesh calls finish.  A malformed/incomplete lattice is ignored and falls back to
     // the canonical generator path rather than changing geometry.
     void SetSharedDensityGrid(const FVoxelSharedDensityGrid* InGrid) { SharedDensityGrid = InGrid; }
+    void SetGridFillExecution(const FVoxelGridFillExecution& InExecution)
+    { GridFillExecution = InExecution; }
 
     // Standalone's interval world is an explicit immutable sampling context. Game builds do not
     // set this pointer and continue through Generator.
@@ -189,6 +211,7 @@ private:
 
     const std::atomic<bool>* ShutdownFlag = nullptr;
     const FVoxelSharedDensityGrid* SharedDensityGrid = nullptr;
+    FVoxelGridFillExecution GridFillExecution;
     const FStandaloneShapeWorld* ShapeWorldContext = nullptr;
     TArray<FVoxelHybridSharpFeatureBox> HybridSharpFeatures;
     bool bUseHybridSharpMesher = false;