cleanup

Author: JimBobSquarePants

src/ImageSharp.Drawing.WebGPU/Shaders/WgslSource/Shared/bump.wgsl

@@ -7,7 +7,6 @@ const STAGE_TILE_ALLOC: u32 = 0x2u;
 const STAGE_FLATTEN: u32 = 0x4u;
 const STAGE_PATH_COUNT: u32 = 0x8u;
 const STAGE_COARSE: u32 = 0x10u;
-const PREVIOUS_RUN: u32 = 0x20u;
 
 // This must be kept in sync with the struct in config.rs in the encoding crate.
 struct BumpAllocators {

src/ImageSharp.Drawing.WebGPU/Shaders/WgslSource/Shared/config.wgsl

@@ -34,7 +34,6 @@ struct Config {
 
     transform_base: u32,
     style_base: u32,
-    cancelled: u32,
 
     // Sizes of bump allocated buffers (in element size units)
     lines_size: u32,

src/ImageSharp.Drawing.WebGPU/Shaders/WgslSource/coarse.wgsl

@@ -171,7 +171,7 @@ fn main(
     // We need to check only prior stages, as if this stage has failed in another workgroup, 
     // we still want to know this workgroup's memory requirement.   
     if local_id.x == 0u {
-        var failed = atomicLoad(&bump.failed) & (STAGE_BINNING | STAGE_TILE_ALLOC | STAGE_FLATTEN | PREVIOUS_RUN);
+        var failed = atomicLoad(&bump.failed) & (STAGE_BINNING | STAGE_TILE_ALLOC | STAGE_FLATTEN);
         if atomicLoad(&bump.seg_counts) > config.seg_counts_size {
             failed |= STAGE_PATH_COUNT;
         }

src/ImageSharp.Drawing.WebGPU/Shaders/WgslSource/flatten.wgsl

@@ -850,9 +850,6 @@ fn main(
     @builtin(global_invocation_id) global_id: vec3<u32>,
     @builtin(local_invocation_id) local_id: vec3<u32>,
 ) {
-    if config.cancelled != 0u {
-        return;
-    }
     let ix = global_id.x;
     pathdata_base = config.pathdata_base;
     bbox = vec4(1e31, 1e31, -1e31, -1e31);

src/ImageSharp.Drawing.WebGPU/Shaders/WgslSource/prepare.wgsl

@@ -12,40 +12,10 @@ var<storage, read_write> bump: BumpAllocators;
 
 @compute @workgroup_size(1)
 fn main() {
-    var should_cancel = false;
-    let previous_failure = atomicLoad(&bump.failed);
-    if (previous_failure & PREVIOUS_RUN) != 0u {
-        atomicStore(&bump.failed, 0u);
-    } else if previous_failure != 0u {
-        if config.binning_size < atomicLoad(&bump.binning) {
-            should_cancel = true;
-        }
-        if config.ptcl_size < atomicLoad(&bump.ptcl) {
-            should_cancel = true;
-        }
-        if config.tiles_size < atomicLoad(&bump.tile) {
-            should_cancel = true;
-        }
-        if config.seg_counts_size < atomicLoad(&bump.seg_counts) {
-            should_cancel = true;
-        }
-        if config.segments_size < atomicLoad(&bump.segments) {
-            should_cancel = true;
-        }
-        if config.lines_size < atomicLoad(&bump.lines) {
-            should_cancel = true;
-        }
-        if config.blend_size < atomicLoad(&bump.blend_spill) {
-            should_cancel = true;
-        }
-        if should_cancel {
-            config.cancelled = 1u;
-            atomicStore(&bump.failed, PREVIOUS_RUN);
-        } else {
-            atomicStore(&bump.failed, 0u);
-        }
-    }
-
+    // Never cancel. Let all stages run so the bump allocators report the true
+    // demand for every buffer in a single pass. The CPU reads back the actuals
+    // and retries once with the correct sizes.
+    atomicStore(&bump.failed, 0u);
     atomicStore(&bump.binning, 0u);
     atomicStore(&bump.ptcl, 0u);
     atomicStore(&bump.tile, 0u);

src/ImageSharp.Drawing.WebGPU/WebGPUSceneDispatch.cs

@@ -162,13 +162,17 @@ public static bool TryCreateStagedScene<TPixel>(
 
         try
         {
+            System.Diagnostics.Stopwatch sw = System.Diagnostics.Stopwatch.StartNew();
             if (!WebGPUSceneEncoder.TryEncode(scene, flushContext.TargetBounds, flushContext.MemoryAllocator, out WebGPUEncodedScene createdScene, out error))
             {
                 flushContext.Dispose();
                 stagedScene = default;
                 return false;
             }
 
+            double encodeMs = sw.Elapsed.TotalMilliseconds;
+            sw.Restart();
+
             encodedScene = createdScene;
             WebGPUSceneConfig config = WebGPUSceneConfig.Create(encodedScene, bumpSizes);
             uint baseColor = 0U;
@@ -202,6 +206,9 @@ public static bool TryCreateStagedScene<TPixel>(
                 return false;
             }
 
+            double resourceMs = sw.Elapsed.TotalMilliseconds;
+            System.IO.File.AppendAllText("bump_debug.log", $"[TIMING] encode={encodeMs:0.000}ms resources={resourceMs:0.000}ms fills={encodedScene.FillCount} paths={encodedScene.PathCount} lines={encodedScene.LineCount}\n");
+
             stagedScene = new WebGPUStagedScene(flushContext, encodedScene, config, resources, segmentChunkingRequired ? bindingLimitFailure : BindingLimitFailure.None);
             error = null;
             return true;
@@ -889,6 +896,8 @@ public static unsafe bool TryRenderStagedScene(
             return TryRenderSegmentChunkedStagedScene(ref stagedScene, ref schedulingArena, out requiresGrowth, out grownBumpSizes, out error);
         }
 
+        var sw = System.Diagnostics.Stopwatch.StartNew();
+
         if (!TryEnsureSchedulingArena(
                 stagedScene.FlushContext,
                 stagedScene.Config.BufferSizes,
@@ -899,31 +908,21 @@ public static unsafe bool TryRenderStagedScene(
             return false;
         }
 
+        double arenaMs = sw.Elapsed.TotalMilliseconds;
+        sw.Restart();
+
         if (!TryDispatchSchedulingStages(ref stagedScene, in schedulingArena, out WebGPUSceneSchedulingResources scheduling, out error))
         {
             return false;
         }
 
-        // Submit scheduling + readback copy so the GPU starts immediately.
-        // While the GPU runs scheduling, the CPU records the fine pass below.
-        if (!TryEnqueueSchedulingStatusReadback(stagedScene.FlushContext, scheduling.BumpBuffer, schedulingArena.ReadbackBuffer, 0, out error) ||
-            !WebGPUDrawingBackend.TrySubmit(stagedScene.FlushContext))
-        {
-            return false;
-        }
+        double schedDispatchMs = sw.Elapsed.TotalMilliseconds;
+        sw.Restart();
 
-        // Record fine + output copy on a new command encoder while GPU runs scheduling.
-        // This CPU work overlaps with GPU scheduling execution.
         WebGPUFlushContext flushContext = stagedScene.FlushContext;
         int targetWidth = encodedScene.TargetSize.Width;
         int targetHeight = encodedScene.TargetSize.Height;
 
-        if (!flushContext.EnsureCommandEncoder())
-        {
-            error = "Failed to create a command encoder for the staged-scene fine pass.";
-            return false;
-        }
-
         if (!WebGPUDrawingBackend.TryCreateCompositionTexture(flushContext, targetWidth, targetHeight, out Texture* outputTexture, out TextureView* outputTextureView, out error))
         {
             return false;
@@ -954,16 +953,24 @@ public static unsafe bool TryRenderStagedScene(
             targetWidth,
             targetHeight);
 
-        // NOW sync-readback the scheduling bumps. By this point the GPU has been
-        // running scheduling while the CPU was recording fine above, so the map
-        // blocks for minimal time (scheduling is usually done or nearly done).
-        if (!TryReadSchedulingStatus(flushContext, schedulingArena.ReadbackBuffer, out GpuSceneBumpAllocators bumpAllocators, out error))
+        double fineDispatchMs = sw.Elapsed.TotalMilliseconds;
+        sw.Restart();
+
+        // Single submit: scheduling + fine + copy + readback all in one.
+        if (!TryEnqueueSchedulingStatusReadback(flushContext, scheduling.BumpBuffer, schedulingArena.ReadbackBuffer, 0, out error) ||
+            !WebGPUDrawingBackend.TrySubmit(flushContext) ||
+            !TryReadSchedulingStatus(flushContext, schedulingArena.ReadbackBuffer, out GpuSceneBumpAllocators bumpAllocators, out error))
         {
             return false;
         }
 
+        double submitReadbackMs = sw.Elapsed.TotalMilliseconds;
+        System.IO.File.AppendAllText("bump_debug.log", $"[RENDER] arena={arenaMs:0.000}ms schedDispatch={schedDispatchMs:0.000}ms fineDispatch={fineDispatchMs:0.000}ms submitReadback={submitReadbackMs:0.000}ms\n");
+
         if (RequiresScratchReallocation(in bumpAllocators, stagedScene.Config.BumpSizes))
         {
+            // Overflow detected. The fine output is garbage but all bump allocators
+            // now report the actual sizes needed. One retry with these sizes will succeed.
             requiresGrowth = true;
             grownBumpSizes = GrowBumpSizes(stagedScene.Config.BumpSizes, in bumpAllocators);
             WebGPUSceneBumpSizes cfg = stagedScene.Config.BumpSizes;
@@ -972,15 +979,7 @@ public static unsafe bool TryRenderStagedScene(
             return false;
         }
 
-        if (string.Equals(Environment.GetEnvironmentVariable("IMAGE_SHARP_WEBGPU_DEBUG_SCHED"), "1", StringComparison.Ordinal))
-        {
-            WebGPUEncodedScene debugScene = stagedScene.EncodedScene;
-            error = $"scene fills={debugScene.FillCount} paths={debugScene.PathCount} lines={debugScene.LineCount} pathtag_bytes={debugScene.PathTagByteCount} pathtag_words={debugScene.PathTagWordCount} drawtags={debugScene.DrawTagCount} drawdata={debugScene.DrawDataWordCount} transforms={debugScene.TransformWordCount} styles={debugScene.StyleWordCount}; sched failed={bumpAllocators.Failed} binning={bumpAllocators.Binning} ptcl={bumpAllocators.Ptcl} tile={bumpAllocators.Tile} seg_counts={bumpAllocators.SegCounts} segments={bumpAllocators.Segments} blend={bumpAllocators.BlendSpill} lines={bumpAllocators.Lines}";
-            return false;
-        }
-
-        // Scheduling passed. Persist the actual GPU usage so the next flush starts
-        // from sizes that are known to work, eliminating retries for similar scenes.
+        // Persist the actual GPU usage so the next flush starts from known-good sizes.
         grownBumpSizes = new WebGPUSceneBumpSizes(
             Math.Max(bumpAllocators.Lines, stagedScene.Config.BumpSizes.Lines),
             Math.Max(bumpAllocators.Binning, stagedScene.Config.BumpSizes.Binning),
@@ -990,8 +989,7 @@ public static unsafe bool TryRenderStagedScene(
             Math.Max(bumpAllocators.BlendSpill, stagedScene.Config.BumpSizes.BlendSpill),
             Math.Max(bumpAllocators.Ptcl, stagedScene.Config.BumpSizes.Ptcl));
 
-        // Submit the pre-recorded fine + copy. Fire-and-forget.
-        return WebGPUDrawingBackend.TrySubmit(flushContext);
+        return true;
     }
 
     /// <summary>