Reduce GPU text memory, page-switch and resize latency

src/ProGPU.Backend/GpuCoverageUpload.cs

@@ -0,0 +1,96 @@
+using Silk.NET.WebGPU;
+
+namespace ProGPU.Backend;
+
+/// <summary>
+/// Records copies from GPU-computed, tightly packed coverage bytes into a filterable R8 atlas.
+/// Rasterization remains GPU-only while avoiding the four-channel storage-texture footprint.
+/// </summary>
+public static unsafe class GpuCoverageUpload
+{
+    public const uint CopyRowAlignment = 256;
+
+    public static uint GetBytesPerRow(uint width)
+    {
+        if (width == 0)
+        {
+            return 0;
+        }
+
+        return checked((width + CopyRowAlignment - 1) & ~(CopyRowAlignment - 1));
+    }
+
+    public static uint GetRequiredBytes(uint width, uint height) =>
+        checked(GetBytesPerRow(width) * height);
+
+    public static void RecordCopy(
+        WgpuContext context,
+        CommandEncoder* encoder,
+        GpuBuffer source,
+        uint sourceOffset,
+        uint bytesPerRow,
+        GpuTexture destination,
+        uint destinationX,
+        uint destinationY,
+        uint width,
+        uint height)
+    {
+        ArgumentNullException.ThrowIfNull(context);
+        ArgumentNullException.ThrowIfNull(source);
+        ArgumentNullException.ThrowIfNull(destination);
+        if (encoder == null)
+        {
+            throw new ArgumentNullException(nameof(encoder));
+        }
+        if (destination.Format != TextureFormat.R8Unorm)
+        {
+            throw new ArgumentException("Coverage copies require an R8Unorm destination.", nameof(destination));
+        }
+        if (!source.Usage.HasFlag(BufferUsage.CopySrc))
+        {
+            throw new ArgumentException("Coverage staging buffers require CopySrc usage.", nameof(source));
+        }
+        if (width == 0 || height == 0)
+        {
+            return;
+        }
+        if (bytesPerRow < width || bytesPerRow % CopyRowAlignment != 0)
+        {
+            throw new ArgumentOutOfRangeException(nameof(bytesPerRow));
+        }
+
+        var copySource = new ImageCopyBuffer
+        {
+            Buffer = source.BufferPtr,
+            Layout = new TextureDataLayout
+            {
+                Offset = sourceOffset,
+                BytesPerRow = bytesPerRow,
+                RowsPerImage = height
+            }
+        };
+        var copyDestination = new ImageCopyTexture
+        {
+            Texture = destination.TexturePtr,
+            MipLevel = 0,
+            Origin = new Origin3D
+            {
+                X = destinationX,
+                Y = destinationY,
+                Z = 0
+            },
+            Aspect = TextureAspect.All
+        };
+        var copySize = new Extent3D
+        {
+            Width = width,
+            Height = height,
+            DepthOrArrayLayers = 1
+        };
+        context.Api.CommandEncoderCopyBufferToTexture(
+            encoder,
+            &copySource,
+            &copyDestination,
+            &copySize);
+    }
+}

src/ProGPU.Backend/GpuTexture.cs

@@ -1201,6 +1201,93 @@ public void CopyBaseLevelFrom(GpuTexture source)
         Generation++;
     }
 
+    /// <summary>
+    /// Copies a top-left base-level region from a compatible texture. This is used
+    /// by demand-grown atlases to preserve resident texels without a CPU readback.
+    /// </summary>
+    public void CopyBaseLevelRegionFrom(GpuTexture source, uint width, uint height)
+    {
+        if (IsDisposed) throw new ObjectDisposedException(nameof(GpuTexture));
+        ArgumentNullException.ThrowIfNull(source);
+        if (source.IsDisposed) throw new ObjectDisposedException(nameof(GpuTexture));
+        if (!ReferenceEquals(source.Context, _context))
+        {
+            throw new ArgumentException(
+                "Source texture must belong to the same WebGPU context.",
+                nameof(source));
+        }
+        if (source.DepthOrArrayLayers != 1 || DepthOrArrayLayers != 1 ||
+            source.Dimension != GpuTextureDimension.Dimension2D ||
+            Dimension != GpuTextureDimension.Dimension2D ||
+            source.Format != Format || source.SampleCount != SampleCount)
+        {
+            throw new ArgumentException(
+                "Source and destination must be compatible single-layer 2D textures.",
+                nameof(source));
+        }
+        if (width == 0 || width > source.Width || width > Width)
+        {
+            throw new ArgumentOutOfRangeException(nameof(width));
+        }
+        if (height == 0 || height > source.Height || height > Height)
+        {
+            throw new ArgumentOutOfRangeException(nameof(height));
+        }
+        if (!source.Usage.HasFlag(TextureUsage.CopySrc))
+        {
+            throw new InvalidOperationException("Source texture was not created with CopySrc usage.");
+        }
+        if (!Usage.HasFlag(TextureUsage.CopyDst))
+        {
+            throw new InvalidOperationException("Destination texture was not created with CopyDst usage.");
+        }
+
+        var encoderDescriptor = new CommandEncoderDescriptor();
+        var encoder = _context.Api.DeviceCreateCommandEncoder(_context.Device, &encoderDescriptor);
+        if (encoder == null)
+        {
+            throw new InvalidOperationException(
+                "Failed to create a command encoder for the regional texture copy.");
+        }
+
+        var copySource = new ImageCopyTexture
+        {
+            Texture = source.TexturePtr,
+            MipLevel = 0,
+            Origin = new Origin3D(),
+            Aspect = GetTextureCopyAspect(source.Format)
+        };
+        var copyDestination = new ImageCopyTexture
+        {
+            Texture = TexturePtr,
+            MipLevel = 0,
+            Origin = new Origin3D(),
+            Aspect = GetTextureCopyAspect(Format)
+        };
+        var copySize = new Extent3D
+        {
+            Width = width,
+            Height = height,
+            DepthOrArrayLayers = 1
+        };
+        _context.Api.CommandEncoderCopyTextureToTexture(
+            encoder,
+            &copySource,
+            &copyDestination,
+            &copySize);
+
+        var commandBufferDescriptor = new CommandBufferDescriptor();
+        var commandBuffer = _context.Api.CommandEncoderFinish(
+            encoder,
+            &commandBufferDescriptor);
+        _context.Api.QueueSubmit(_context.Queue, 1, &commandBuffer);
+        _context.Api.CommandBufferRelease(commandBuffer);
+        _context.Api.CommandEncoderRelease(encoder);
+
+        AlphaMode = source.AlphaMode;
+        Generation++;
+    }
+
     private void EnsurePbgra32CompatibleFormat()
     {
         if (Format is not (TextureFormat.Bgra8Unorm or TextureFormat.Bgra8UnormSrgb))

src/ProGPU.Backend/IWebGpuApi.cs

@@ -29,6 +29,7 @@ public unsafe interface IWebGpuApi
     ComputePassEncoder* CommandEncoderBeginComputePass(CommandEncoder* commandEncoder, ComputePassDescriptor* descriptor);
     RenderPassEncoder* CommandEncoderBeginRenderPass(CommandEncoder* commandEncoder, RenderPassDescriptor* descriptor);
     void CommandEncoderCopyBufferToBuffer(CommandEncoder* commandEncoder, WgpuBuffer* source, ulong sourceOffset, WgpuBuffer* destination, ulong destinationOffset, ulong size);
+    void CommandEncoderCopyBufferToTexture(CommandEncoder* commandEncoder, ImageCopyBuffer* source, ImageCopyTexture* destination, Extent3D* copySize);
     void CommandEncoderCopyTextureToBuffer(CommandEncoder* commandEncoder, ImageCopyTexture* source, ImageCopyBuffer* destination, Extent3D* copySize);
     void CommandEncoderCopyTextureToTexture(CommandEncoder* commandEncoder, ImageCopyTexture* source, ImageCopyTexture* destination, Extent3D* copySize);
     CommandBuffer* CommandEncoderFinish(CommandEncoder* commandEncoder, CommandBufferDescriptor* descriptor);
@@ -118,6 +119,7 @@ private static void CompleteBufferMap(BufferMapAsyncStatus status, void* userDat
     public ComputePassEncoder* CommandEncoderBeginComputePass(CommandEncoder* e, ComputePassDescriptor* x) => api.CommandEncoderBeginComputePass(e, x);
     public RenderPassEncoder* CommandEncoderBeginRenderPass(CommandEncoder* e, RenderPassDescriptor* x) => api.CommandEncoderBeginRenderPass(e, x);
     public void CommandEncoderCopyBufferToBuffer(CommandEncoder* e, WgpuBuffer* s, ulong so, WgpuBuffer* d, ulong @do, ulong z) => api.CommandEncoderCopyBufferToBuffer(e, s, so, d, @do, z);
+    public void CommandEncoderCopyBufferToTexture(CommandEncoder* e, ImageCopyBuffer* s, ImageCopyTexture* d, Extent3D* z) => api.CommandEncoderCopyBufferToTexture(e, s, d, z);
     public void CommandEncoderCopyTextureToBuffer(CommandEncoder* e, ImageCopyTexture* s, ImageCopyBuffer* d, Extent3D* z) => api.CommandEncoderCopyTextureToBuffer(e, s, d, z);
     public void CommandEncoderCopyTextureToTexture(CommandEncoder* e, ImageCopyTexture* s, ImageCopyTexture* d, Extent3D* z) => api.CommandEncoderCopyTextureToTexture(e, s, d, z);
     public CommandBuffer* CommandEncoderFinish(CommandEncoder* e, CommandBufferDescriptor* x) => api.CommandEncoderFinish(e, x);

src/ProGPU.Backend/Shaders/GlyphRasterizer.wgsl

@@ -1,13 +1,13 @@
 // Algorithm: Rasterize glyph coverage with 8x8 supersampling, sharing analytic line, quadratic, and cubic winding intersections across each eight-sample row.
 // Time complexity: O(R*S + A) per texel for R=8 sample rows, A=64 anti-aliasing samples, and S outline segments.
-// Space complexity: O(R) private winding storage plus O(S) read-only segment bandwidth and one rgba8unorm output write per texel.
+// Space complexity: O(R) private winding storage plus O(S) read-only segment bandwidth and one packed u32 output write per four R8 texels.
 struct GlyphUniforms {
     xStart: f32,
     yStart: f32,
     scale: f32,
     glyphIndex: u32,
-    atlasX: u32,
-    atlasY: u32,
+    outputOffsetWords: u32,
+    outputRowWords: u32,
     width: u32,
     height: u32,
     subpixelX: f32,
@@ -41,7 +41,7 @@ struct Segment {
 @group(0) @binding(0) var<uniform> uniforms: GlyphUniforms;
 @group(0) @binding(1) var<storage, read> glyphRecords: array<GlyphRecord>;
 @group(0) @binding(2) var<storage, read> segments: array<Segment>;
-@group(0) @binding(3) var atlasTexture: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(3) var<storage, read_write> coverageOutput: array<u32>;
 
 
 fn solve_quadratic(a: f32, b: f32, c: f32, roots: ptr<function, array<f32, 2>>, root_count: ptr<function, u32>) {
@@ -269,15 +269,7 @@ fn accumulate_winding_row(
     }
 }
 
-@compute @workgroup_size(16, 16)
-fn cs_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
-    let x = global_id.x;
-    let y = global_id.y;
-
-    if (x >= uniforms.width || y >= uniforms.height) {
-        return;
-    }
-
+fn glyph_coverage_byte(x: u32, y: u32) -> u32 {
     let glyphIndex = uniforms.glyphIndex;
     let record = glyphRecords[glyphIndex];
 
@@ -305,7 +297,27 @@ fn cs_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
         }
     }
 
-    let coverage = f32(covered_samples) * 0.015625;
-    let writeCoord = vec2<u32>(uniforms.atlasX + x, uniforms.atlasY + y);
-    textureStore(atlasTexture, writeCoord, vec4<f32>(coverage, 0.0, 0.0, 0.0));
+    return min(255u, u32(round(f32(covered_samples) * 3.984375)));
+}
+
+// Four adjacent pixels are packed by one invocation so the storage buffer has
+// the exact byte layout required by WebGPU copyBufferToTexture for R8Unorm.
+@compute @workgroup_size(16, 16)
+fn cs_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    let wordX = global_id.x;
+    let y = global_id.y;
+    let firstX = wordX * 4u;
+    if (firstX >= uniforms.width || y >= uniforms.height) {
+        return;
+    }
+
+    var packed = 0u;
+    for (var lane = 0u; lane < 4u; lane = lane + 1u) {
+        let x = firstX + lane;
+        if (x < uniforms.width) {
+            packed = packed | (glyph_coverage_byte(x, y) << (lane * 8u));
+        }
+    }
+
+    coverageOutput[uniforms.outputOffsetWords + y * uniforms.outputRowWords + wordX] = packed;
 }

src/ProGPU.Backend/Shaders/PathRasterizer.wgsl

@@ -1,14 +1,14 @@
 // Algorithm: Rasterize arbitrary path coverage by supersampling each atlas texel and applying analytic non-zero or even-odd winding tests.
 // Time complexity: O(A*S) per texel for A anti-aliasing samples and S path segments.
-// Space complexity: O(1) local storage plus O(S) read-only segment bandwidth.
+// Space complexity: O(1) local storage plus O(S) read-only segment bandwidth and one packed u32 output write per four R8 texels.
 struct PathUniforms {
     xStart: f32,
     yStart: f32,
     scaleX: f32,
     scaleY: f32,
     pathIndex: u32,
-    atlasX: u32,
-    atlasY: u32,
+    outputOffsetWords: u32,
+    outputRowWords: u32,
     width: u32,
     height: u32,
     sampleGrid: u32,
@@ -41,7 +41,7 @@ struct Segment {
 @group(0) @binding(0) var<storage, read> pathUniforms: array<PathUniforms>;
 @group(0) @binding(1) var<storage, read> pathRecords: array<PathRecord>;
 @group(0) @binding(2) var<storage, read> segments: array<Segment>;
-@group(0) @binding(3) var atlasTexture: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(3) var<storage, read_write> coverageOutput: array<u32>;
 
 
 fn solve_quadratic(a: f32, b: f32, c: f32, roots: ptr<function, array<f32, 2>>, root_count: ptr<function, u32>) {
@@ -341,16 +341,7 @@ fn count_row_coverage(
     return covered;
 }
 
-@compute @workgroup_size(16, 16)
-fn cs_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
-    let uniforms = pathUniforms[global_id.z];
-    let x = global_id.x;
-    let y = global_id.y;
-
-    if (x >= uniforms.width || y >= uniforms.height) {
-        return;
-    }
-
+fn path_coverage_byte(x: u32, y: u32, uniforms: PathUniforms) -> u32 {
     let pathIndex = uniforms.pathIndex;
     let record = pathRecords[pathIndex];
 
@@ -369,8 +360,28 @@ fn cs_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
             uniforms.scaleX,
             record);
     }
-    let coverage = f32(coveredSamples) * sampleWeight;
+    return min(255u, u32(round(f32(coveredSamples) * sampleWeight * 255.0)));
+}
+
+// Four adjacent pixels are packed by one invocation so the storage buffer has
+// the exact byte layout required by WebGPU copyBufferToTexture for R8Unorm.
+@compute @workgroup_size(16, 16)
+fn cs_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    let uniforms = pathUniforms[global_id.z];
+    let wordX = global_id.x;
+    let y = global_id.y;
+    let firstX = wordX * 4u;
+    if (firstX >= uniforms.width || y >= uniforms.height) {
+        return;
+    }
+
+    var packed = 0u;
+    for (var lane = 0u; lane < 4u; lane = lane + 1u) {
+        let x = firstX + lane;
+        if (x < uniforms.width) {
+            packed = packed | (path_coverage_byte(x, y, uniforms) << (lane * 8u));
+        }
+    }
 
-    let writeCoord = vec2<u32>(uniforms.atlasX + x, uniforms.atlasY + y);
-    textureStore(atlasTexture, writeCoord, vec4<f32>(coverage, 0.0, 0.0, 0.0));
+    coverageOutput[uniforms.outputOffsetWords + y * uniforms.outputRowWords + wordX] = packed;
 }