From 781d95ae75382355852efa0dc90fa1cb9e067374 Mon Sep 17 00:00:00 2001 From: Jorijn van der Graaf Date: Sat, 31 Oct 2026 20:22:12 +0100 Subject: [PATCH] fix(ui): decode the font atlas's SDF in atlas pixels, so text edges ramp over one screen pixel Shaders - ui-text and ui-fused turned a sample of the R8_UNORM atlas into a distance with (128/255 - sdf) * 32. stb_truetype writes kOnEdgeValue + kPixelDistScale * distance into each byte, so an atlas pixel is kPixelDistScale / 255 of a sample; 32 made every distance 1024/255 (about 4) times too large, and the coverage ramp, meant to span one screen pixel, spanned a quarter of one: nearly every edge pixel came out fully on or fully off, and text looked pixelated at every size. On an RTX 4090, UIFallbackParity's 22 px label had 98 partly covered pixels where stb_truetype's own rasterizer has 415. Wrong since the v3 UI (1f56973), and copied into the fused kernel (574242d). - The decode now lives once, in ui-shared.glsl: UI_SDF_ON_EDGE and UI_SDF_PIXEL_DIST_SCALE are FontAtlas's two constants in the encoder's own units, uiSdfDistance decodes a sample, uiGlyphRamp works out a glyph's ramp when it is loaded and uiGlyphCoverage the coverage per pixel. Both kernels call them, and a copy of either should; the UI_ prefix keeps them clear of a copied kernel's own names. - stb_truetype truncates, so a byte b stands for [b, b + 1) of the encoded value and the decode takes its middle: decoding b pulled every outline 1/64 atlas pixel inwards. - The ramp is capped at what a byte holds, about 4 atlas pixels each way of the outline. With the corrected decode a wider one left every glyph quad's padding translucent below 4 px. - The UV span is taken with abs(): a glyph turned by swapping its UVs got a negative scale, a ramp of 0.0001 and hard edges. - A glyph scaled differently along x and y ramps along the distance's gradient, from two one-texel differences. With the corrected decode, one width for both axes, the more minified one's, would leave a row stretched three times with stems three pixels soft. Glyphs scaled alike, as ShapeText's are, keep one multiply-add per pixel, the ramp's reciprocal worked out with the glyph. - The atlas size is queried rather than FontAtlas::kAtlasSize hard-coded as 1024. WebGPU - The WGSL text kernel decoded the distance right but ramped over one atlas pixel whatever the size: half a screen pixel at 16 px, two at 64 px. It now repeats ui-shared.glsl's decode, cap, abs and gradient path, and renders TextAntialiasing's scene as Vulkan does (run through Deno's WebGPU: 1.02, 1.07 and 1.09 times the reference's edge pixels, where the old kernel gave 0.55, 2.19 and 0.89). FontAtlas - kOnEdgeValue and kPixelDistScale say the kernels carry copies of them. Tests - TextAntialiasing draws one string at 14 px, 64 px and 14 px stretched three times, a 'g' turned by its UVs and empty atlas texels shrunk to 12.8 atlas pixels per screen pixel, through DispatchText and again through DispatchFused, and holds them to stb_truetype's antialiased rasterizer, to the mirror image and to nothing. On master the two plain rows have 0.31 and 0.21 times the reference's edge pixels; with the decode fixed alone the stretched row had 1.62 times, the turned 'g' was 127 codes off its mirror and the empty texels drew 48. Now: 1.02, 1.07 and 1.09 times, 0 and 0, on the 4090, on llvmpipe and on descriptor sets. It skips a swapchain that isn't 8-bit, which the capture can't read, and captures the first frame. - FontAtlasDecode reads ui-shared.glsl and dom-webgpu.js and holds their copies of the two constants to FontAtlas's. No GPU. Docs - AGENTS.consumer.md has a custom text shader decode through the helpers and gives the encoding; README says so in a line; AGENTS.md lists the tests and where the decode lives. Co-Authored-By: Claude Opus 5.5 --- AGENTS.consumer.md | 3 +- AGENTS.md | 7 +- README.md | 6 +- additional/dom-webgpu.js | 37 +- interfaces/Crafter.Graphics-FontAtlas.cppm | 5 + project.cpp | 22 + shaders/ui-fused.comp.glsl | 36 +- shaders/ui-shared.glsl | 87 ++++ shaders/ui-text.comp.glsl | 40 +- tests/FontAtlasDecode/main.cpp | 101 +++++ tests/TextAntialiasing/main.cpp | 486 +++++++++++++++++++++ 11 files changed, 767 insertions(+), 63 deletions(-) create mode 100644 tests/FontAtlasDecode/main.cpp create mode 100644 tests/TextAntialiasing/main.cpp diff --git a/AGENTS.consumer.md b/AGENTS.consumer.md index 175e3f6..4bc6d3c 100644 --- a/AGENTS.consumer.md +++ b/AGENTS.consumer.md @@ -83,7 +83,7 @@ Per frame, in an `EventListener` on `ui.onBuild` (fired from the pa - The item structs are `std430` PODs mirrored byte for byte in `shaders/ui-shared.glsl`, and frozen: `QuadItem` (64 bytes: `x, y, w, h`, `r, g, b, a`, corner radii `cTL, cTR, cBR, cBL`, `outline, oR, oG, oB`), `CircleItem` (48: `cx, cy, radius, _p0`, colour, outline), `ImageItem` (64: rect, `u0, v0, u1, v1`, tint, `texSlot, sampSlot, _p1, _p2`), `GlyphItem` (48: rect, atlas uv, colour). Colours are straight alpha, blended source-over; coordinates are window pixels. - Cost: every kernel dispatches one workgroup per 8x8 screen tile and each pixel walks every item (behind a per-tile cull), so a dispatch is `O(W*H*N)`, fine up to a few hundred items at 1080p. Splitting one buffer into several dispatches does not help. Thousands of items want a different shader (tile binning), which is Tier 1. -- Text: `Font font("font.ttf")` (TrueType through stb_truetype); `font.GetLineWidth(text, px)`, `LineHeight(px)`, `AscentPx(px)`, `NearestCursorByte(text, px, x)`. The atlas is a 1024 by 1024 SDF atlas at a 32 px base size, filled lazily per codepoint and re-uploaded by dirty rectangle; glyphs scale from it, so very large text goes soft. `ui.ShapeText(font, px, x, baselineY, utf8, color, out, capacity, &advance, TextAlign::Left)` writes `GlyphItem`s: single line, no kerning, no wrapping, a shaped-run cache behind it. Once the atlas is full, a glyph that finds no free room draws blank for one frame but keeps its advance, and gets a cell when the next frame begins, taken from the glyphs drawn least recently, never one drawn in the last `Device::framesInFlight + 1` frames. So a `GlyphItem` holds for the frame it was shaped for: shape text every frame (the Tier 3 components do, and a cached run costs a copy) rather than drawing a glyph buffer filled in an earlier frame. `atlas.EnsureAndGet(font, codepoint)` returns `const Glyph&` and never fails; its UVs hold only when `Resident()`, and a `Lookup` pointer only until the atlas's next `EnsureAndGet`, `Touch` or `Update`. A Tier 1 kernel that samples `FontAtlasImageSlot()` through `ui.Dispatch` sees the frame's new glyphs: every dispatch uploads the atlas first. +- Text: `Font font("font.ttf")` (TrueType through stb_truetype); `font.GetLineWidth(text, px)`, `LineHeight(px)`, `AscentPx(px)`, `NearestCursorByte(text, px, x)`. The atlas is a 1024 by 1024 SDF atlas at a 32 px base size, filled lazily per codepoint and re-uploaded by dirty rectangle. Glyphs scale from it with the edge ramped over one screen pixel at any size and aspect, on both backends; far above 32 px corners round off and fine detail goes, and below about 4 px the ramp narrows and thin strokes bolden. `ui.ShapeText(font, px, x, baselineY, utf8, color, out, capacity, &advance, TextAlign::Left)` writes `GlyphItem`s: single line, no kerning, no wrapping, a shaped-run cache behind it. Once the atlas is full, a glyph that finds no free room draws blank for one frame but keeps its advance, and gets a cell when the next frame begins, taken from the glyphs drawn least recently, never one drawn in the last `Device::framesInFlight + 1` frames. So a `GlyphItem` holds for the frame it was shaped for: shape text every frame (the Tier 3 components do, and a cached run costs a copy) rather than drawing a glyph buffer filled in an earlier frame. `atlas.EnsureAndGet(font, codepoint)` returns `const Glyph&` and never fails; its UVs hold only when `Resident()`, and a `Lookup` pointer only until the atlas's next `EnsureAndGet`, `Touch` or `Update`. A Tier 1 kernel that samples `FontAtlasImageSlot()` through `ui.Dispatch` sees the frame's new glyphs: every dispatch uploads the atlas first. - Tier 3: fill a `UIBuffer` (`quads`, `quadCount`, `quadCap`; `glyphs`, `glyphCount`, `glyphCap`; `images`, `imageCount`, `imageCap`; `atlas`; `renderer`; a null pointer makes the components that need it no-ops) and call `DrawButton(buf, rect, label, hovered, pressed, font, px, ButtonColors)`, `DrawCheckbox(buf, rect, checked, hovered, CheckboxColors)`, `DrawSlider(buf, rect, t01, dragging, SliderColors)`, `DrawProgressBar(buf, rect, t01, ProgressColors)`, `DrawText(buf, text, x, baselineY, font, px, color, align)` (returns the advance) and `DrawImage(buf, rect, textureSlot, samplerSlot, tint, uv)`. Colours are small `*Colors` aggregates you hold; there is no `Theme`. - `InputField` is the model for a stateful component: a POD (`value`, `type`, `focused`, `cursorPos`), free functions you call from your own listeners (`InputField_OnText(field, utf8)` from `onTextInput`, `InputField_OnKey(field, code)` from `onRawKeyDown`, `InputField_HitTestCursor(field, rect, clickX, font, px, colors)`), `InputField_TryParse(field)`, and `DrawInputField(buf, field, rect, font, px, colors, caretVisible)`; the blink is yours. - Images in the UI: `auto tex = ui.RegisterImage(image, VK_FORMAT_R8G8B8A8_UNORM);` (an `ImageSlot`; the layout defaults to `SHADER_READ_ONLY_OPTIMAL`), `auto samp = ui.RegisterLinearClampSampler();`, then both go into `ImageItem::texSlot` and `sampSlot`. `ui.OutImageSlot()`, `FontAtlasImageSlot()` and `FontAtlasSamplerSlot()` expose the renderer's own slots. @@ -110,6 +110,7 @@ void main() { - Declare bindless resources with `UI_HEAP_IMAGES`, `UI_HEAP_TEXTURES`, `UI_HEAP_SAMPLERS` and `UI_HEAP_BUFFERS`, never a literal `descriptor_heap` or a set and binding, and register the source twice in `project.cpp`: once plain, once with `sets.defines = {"CRAFTER_GRAPHICS_DESCRIPTOR_SETS"}; sets.variant = "sets";`, which yields `.spv` and `.sets.spv`. `ComputeShader::Load(".spv", heap)` picks the build the device needs and throws naming a missing file. `examples/CustomShader/project.cpp` has the lines. - `UIDispatchHeader` (48 bytes; in C++ `outImage, itemBuffer, surfaceWidth, surfaceHeight, clipX, clipY, clipW, clipH, itemCount, frameIdx, flags, _pad`, and in GLSL the same bytes as `uint outImage, itemBuffer; uvec2 surfaceSize; vec4 clipRectPx; uint itemCount, frameIdx, flags, _pad`) must be the first member of the push block; `ui.FillHeader(slot, count, clipRectPx, flags)` fills it. The high bit of `flags` (`kUIFlagClip`) is the renderer's, the rest are yours. Keep the push block a multiple of 4 bytes and within 256; the descriptor-set path throws past the device's `maxPushConstantsSize`. - `ui.Dispatch(cmd, shader, &pc, sizeof(pc), ceil(W / 8), ceil(H / 8))` binds, pushes and dispatches, with a `SHADER_WRITE -> SHADER_READ | SHADER_WRITE` barrier on the swapchain image before every dispatch but the frame's first, so a custom shader sees what the dispatch before it drew. `ComputeShader::Dispatch(cmd, push, bytes, gx, gy, gz, tlasAddress = 0)` is the same without the barrier; a compute shader that ray-queries the TLAS passes `RenderingElement3D::tlases[frameIdx].address`, for the NVIDIA workaround below. +- A shader that draws text from the font atlas (`ui.FontAtlasImageSlot()` and `FontAtlasSamplerSlot()`), a copy of `ui-text.comp.glsl` among them, decodes it with `ui-shared.glsl`'s helpers rather than arithmetic of its own: `uiGlyphAtlasSize(texSlot, samplerSlot)`, then `uiGlyphRamp(item.uv, 1.0 / item.rect.zw, atlasSize)` once per glyph and `uiGlyphCoverage(texSlot, samplerSlot, uv, ramp, atlasSize)` per pixel, with dynamically uniform slots. A texel holds stb_truetype's `UI_SDF_ON_EDGE + UI_SDF_PIXEL_DIST_SCALE * d` (128 + 32 d), `d` the distance from the outline in atlas pixels, positive inside, truncated to a byte; `uiSdfDistance(sample)` gives `d` back. Multiplying a sample by 32 instead of dividing by 32/255 gives a ramp a quarter of a pixel wide, and pixelated text. - Keep the one-workgroup-per-screen-tile shape for anything that composites. A workgroup per item races on overlapping pixels: Vulkan orders nothing between the workgroups of one dispatch. - A `ComputeShader` can also run outside the UI pass, dispatched from `onUpdate` into the frame's command buffer with barriers of your own. diff --git a/AGENTS.md b/AGENTS.md index 7393157..999b17f 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -70,9 +70,9 @@ Tests are declared in `project.cpp`, not discovered, through the `LibraryTest(na A test is an `int main()`: exit 0 passes, nonzero fails, 77 skips (for a missing capability, never for a failed check). Each runs in its own process, in parallel, with a 60 s default timeout, and the output of a failed, crashed or timed-out test is in `build/test-logs/.log`; a later pass does not delete the log, so check its timestamp. What the tests need, in three groups: -- Pure logic, no GPU: they stage `Device`'s statics or feed synthetic data to the decision functions that were factored out for them (`Device::ResolveFeatures`, `ChooseDevice`, `CheckShaderCapabilities`, `GetMemoryType`, `PreferDirectDeviceWrite`, `PipelineCacheDataCompatible`, `AlignMappedFlushRange`, `BuildMipChainFinalBarriers`, `SwapchainStageUnion`, `BuildTlasWriteBarrier`, `UIRenderer::ClipFlags`, `Window::SwapchainMinImageCount`, `WorkaroundNvidiaAS::Patch`, `DescriptorHeapVulkan::ComputeLayout`, `FontAtlas::Packer`, the font and input-field helpers). `PushConstantRewrite` and `UIFusedShader` compile GLSL with glslang and check the SPIR-V with spirv-val at runtime, and `ExampleShaders` compiles every GLSL shader under `examples/` the way crafter-build would, the descriptor-set build included, so a shader contract change that breaks an example fails here without building the example; `MouseScroll` and `PointerEnterLeave` drive the Wayland listener callbacks directly, and `TitleWayland` reads the requests a window sends off a `wl_display` connected to a socketpair. +- Pure logic, no GPU: they stage `Device`'s statics or feed synthetic data to the decision functions that were factored out for them (`Device::ResolveFeatures`, `ChooseDevice`, `CheckShaderCapabilities`, `GetMemoryType`, `PreferDirectDeviceWrite`, `PipelineCacheDataCompatible`, `AlignMappedFlushRange`, `BuildMipChainFinalBarriers`, `SwapchainStageUnion`, `BuildTlasWriteBarrier`, `UIRenderer::ClipFlags`, `Window::SwapchainMinImageCount`, `WorkaroundNvidiaAS::Patch`, `DescriptorHeapVulkan::ComputeLayout`, `FontAtlas::Packer`, the font and input-field helpers). `PushConstantRewrite` and `UIFusedShader` compile GLSL with glslang and check the SPIR-V with spirv-val at runtime, and `ExampleShaders` compiles every GLSL shader under `examples/` the way crafter-build would, the descriptor-set build included, so a shader contract change that breaks an example fails here without building the example; `FontAtlasDecode` reads `ui-shared.glsl` and `dom-webgpu.js` and holds their copies of the SDF encoding constants to `FontAtlas`'s; `MouseScroll` and `PointerEnterLeave` drive the Wayland listener callbacks directly, and `TitleWayland` reads the requests a window sends off a `wl_display` connected to a socketpair. - A headless device: `Device::Initialize()` and no `Window` (`HeapSlotReuse`, `DeferredDeletion`, `MeshLifetime`, `ImageStagingRelease`, `BLASBuildOptions`, the `TLAS*` tests, ...). Headless still means a Wayland display on Linux. Most of the RT ones check `Device::rayTracingSupported` and return 77 without it; `BLASBuildOptions`, `TLASInstanceDirtyTracking` and `MeshDecompressStagingRelease` do not, and crash on a GPU without ray tracing (`Device::RequireRayTracing` throws out of `main`). `DeviceSelection` skips with one GPU. -- A window and a live compositor: `FrameLoopSync`, `SetupCmdBufferReuse`, `UIFallbackParity`, `SwapchainTransferDst`, `WindowUpdateClock`, `HeapASRead`, `HeapDynamicIndexHit`, `RTPassPushData`, `WindowOptionalGlobals` (hides xdg-decoration, viewporter and fractional scale from the window in turn), `SwapchainImageCount` (every present mode, and swapchains with more images than frames in flight, in child processes on llvmpipe where the host's GPU never makes one), and `FeatureRequests` for its retry case. They render, read the frame back with `Window::CaptureNextFrame` (a PNG, decoded with a private `STB_IMAGE_STATIC` copy of stb_image) or a transfer pass of their own, and check pixels. +- A window and a live compositor: `FrameLoopSync`, `SetupCmdBufferReuse`, `UIFallbackParity`, `SwapchainTransferDst`, `WindowUpdateClock`, `HeapASRead`, `HeapDynamicIndexHit`, `RTPassPushData`, `WindowOptionalGlobals` (hides xdg-decoration, viewporter and fractional scale from the window in turn), `SwapchainImageCount` (every present mode, and swapchains with more images than frames in flight, in child processes on llvmpipe where the host's GPU never makes one), `TextAntialiasing` (text through `DispatchText` and `DispatchFused` against stb_truetype's own antialiased rasterizer, at 14 and 64 px and stretched, plus a glyph mirrored by its UVs and empty texels shrunk below 4 px), and `FeatureRequests` for its retry case. They render, read the frame back with `Window::CaptureNextFrame` (a PNG, decoded with a private `STB_IMAGE_STATIC` copy of stb_image) or a transfer pass of their own, and check pixels. Conventions the existing tests follow: @@ -83,7 +83,7 @@ Conventions the existing tests follow: - Files go under `fs::temp_directory_path()`, named per test. - `HeapASRead --no-workaround` is the manual probe for whether a new NVIDIA driver still has the descriptor-heap AS fault; run it after a driver update and update the README's status note. -A bug fix ships with a test that fails before the fix and passes after it. A change to a UI kernel keeps `UIFallbackParity` green: the heap and descriptor-set builds of one source must stay pixel-identical. A category drawn fused must also match the same category drawn standalone, but no test compares those two (`UIFusedShader` only pins the push-constant layout): check it by eye, or add the comparison. +A bug fix ships with a test that fails before the fix and passes after it. A change to a UI kernel keeps `UIFallbackParity` green: the heap and descriptor-set builds of one source must stay pixel-identical. A category drawn fused must also match the same category drawn standalone. `TextAntialiasing` compares the two for text, and keeps the text's edge ramp at one screen pixel; for quads, circles and images no test does (`UIFusedShader` only pins the push-constant layout): check those by eye, or add the comparison. Cross targets: `crafter-build test --target=x86_64-w64-mingw32` needs `wine` on the sandbox's `PATH`, which the Flatpak does not have; a MinGW import library for `vulkan-1`, which the toolchain does not ship; `-luuid`, which the Win32 link flags lack (#184); and a definition of `__libc_single_threaded`, which the MinGW build of Crafter.Asset references. The Win32 tests have been run with an import library made by `llvm-dlltool` from the `vk*` symbols `llvm-nm -u` lists for the MinGW `libCrafter.Graphics.a`, and an archive defining `char __libc_single_threaded = 0;`, both under the gitignored `build/`. A temporary edit of `project.cpp` added `-L` to that directory, both libraries and `-luuid` to the Win32 tests' link flags and cleared their `requires_`; the tests were built with `--runner=cmd:true`, and each exe ran from its bin directory under the host's wine in a scratch `WINEPREFIX`. For an aarch64 device, `crafter-build test --target=...` tries qemu, which has no Vulkan: build with the sysroot, copy the test's whole bin directory (binary, `*.spv`, `*.sets.spv`, fonts) to the device and run it there, unlocked and with the window visible, since a locked phone's compositor sends no frame callbacks and the FIFO present then blocks forever. @@ -150,6 +150,7 @@ These were decided deliberately. Do not undo them without the maintainer's say-s - `BuildTlasWriteBarrier`'s source access mask carries no read bit on purpose: synchronization validation 1.4.363 otherwise reports two builds in one command buffer as a write-after-write hazard. - NVIDIA, as of driver `615.78.08`: reading a TLAS through the heap still faults (worked around); dynamic descriptor indexing in hit shaders works there and faulted on `610.43.02`, so the Sponza pattern (a spec-constant slot, a dynamic index within the resource) stays the portable one. `HeapDynamicIndexHit` skips on the older drivers. - The README links `WAVEFRONT-DESIGN.md`, which was removed with the stale docs (`15ce47f`); the wavefront tracer is described in the README's WebGPU section and in `additional/dom-webgpu.js`. +- The text kernels decode `FontAtlas`'s SDF in one place, `ui-shared.glsl` (`UI_SDF_ON_EDGE`, `UI_SDF_PIXEL_DIST_SCALE`, `uiGlyphRamp`, `uiGlyphCoverage`), and the WGSL text kernel in `dom-webgpu.js` repeats it; `FontAtlasDecode` holds both copies of the constants to `FontAtlas`'s, and nothing renders the WGSL, so check a change to it by hand (`examples/HelloUI` built with `--target=wasm32-wasip1`, in a browser, or the kernel run through Deno's WebGPU). The WGSL lives in `String.raw` literals: a backtick in a WGSL comment ends the literal and breaks the whole backend, which only `node --check additional/dom-webgpu.js` notices. - Crafter.Math's misspelled identifiers (`Quanternion*`, `MulitplyAdd`, `CollumSize`) are its API; do not correct them here. - Dependency notes: crafter-build 3.0.0 or newer keeps a section at the end of this file holding the `AGENTS.consumer.md` of every dependency that ships one (2.2.0 leaves the file alone), rewritten on every `crafter-build` and `crafter-build test`. As of that date Crafter.Math's notes are on its `master`, Crafter.Asset's only on its `claude/glb-loader` branch, and Crafter.Event has none, so a plain build adds Math's and a `--local` build whatever the sibling checkouts hold. Edits inside the section's markers are lost on the next build; a change to those notes belongs in their repositories. Commit the section when it changes. diff --git a/README.md b/README.md index a9be05e..58490b7 100644 --- a/README.md +++ b/README.md @@ -185,7 +185,11 @@ shaded through an intersection shader with an any-hit cut-out. shape text every frame, as the Tier 3 components do (`ShapeText`'s run cache makes that a copy), into item buffers of the frame's own, one per frame in flight (`UIBuildArgs::frameIdx`, as in HelloUI). - Backend-agnostic. + Backend-agnostic. The text kernels decode it through + [shaders/ui-shared.glsl](shaders/ui-shared.glsl)'s `uiGlyphRamp` and + `uiGlyphCoverage`, which a custom text shader should use too (the + WebGPU kernel repeats them in WGSL); a glyph's edge ramps over one + screen pixel at any size and aspect. - **Mesh / RenderingElement3D / Animation** — BLAS/TLAS construction and 3D scene plumbing. Vulkan calls `vkCmdBuildAccelerationStructures`; WebGPU registers BLAS data (verts, idx, BVH nodes, primRemap, optional diff --git a/additional/dom-webgpu.js b/additional/dom-webgpu.js index 5bee542..6db1be3 100644 --- a/additional/dom-webgpu.js +++ b/additional/dom-webgpu.js @@ -439,6 +439,38 @@ struct GlyphItem { @group(3) @binding(0) var atlasTex : texture_2d; @group(3) @binding(1) var atlasSampler : sampler; +// The SDF decode of ui-shared.glsl (uiSdfDistance, uiGlyphRamp and +// uiGlyphCoverage; the reasons are there), worked out per pixel since this +// kernel has no per-glyph shared-memory pass. The FontAtlasDecode test holds +// the two constants to FontAtlas's. +const UI_SDF_ON_EDGE : f32 = 128.0; // FontAtlas::kOnEdgeValue +const UI_SDF_PIXEL_DIST_SCALE : f32 = 32.0; // FontAtlas::kPixelDistScale +const UI_SDF_MAX_BAND : f32 = 2.0 * min(UI_SDF_ON_EDGE - 0.5, 255.5 - UI_SDF_ON_EDGE) / UI_SDF_PIXEL_DIST_SCALE; + +fn uiSdfDistance(sdf: f32) -> f32 { + return (UI_SDF_ON_EDGE - (sdf * 255.0 + 0.5)) / UI_SDF_PIXEL_DIST_SCALE; +} + +// The coverage at uv of the glyph drawn from glyphUv (u0, v0, u1, v1) over a +// rect of rectSize. +fn uiGlyphCoverage(uv: vec2, glyphUv: vec4, rectSize: vec2) -> f32 { + let atlasSize = vec2(textureDimensions(atlasTex)); + let atlasPerScreen = abs(glyphUv.zw - glyphUv.xy) * atlasSize / rectSize; + let sdf = textureSampleLevel(atlasTex, atlasSampler, uv, 0.0).r; + let lo = min(atlasPerScreen.x, atlasPerScreen.y); + let hi = max(atlasPerScreen.x, atlasPerScreen.y); + var band = hi; + if (hi > lo * 1.01) { + let texel = 1.0 / atlasSize; + let gradient = vec2( + textureSampleLevel(atlasTex, atlasSampler, uv + vec2(texel.x, 0.0), 0.0).r, + textureSampleLevel(atlasTex, atlasSampler, uv + vec2(0.0, texel.y), 0.0).r) - sdf; + let len = length(gradient); + if (len > 0.0) { band = length(gradient / len * atlasPerScreen); } + } + return clamp(0.5 - uiSdfDistance(sdf) / clamp(band, 0.0001, UI_SDF_MAX_BAND), 0.0, 1.0); +} + @compute @workgroup_size(8, 8, 1) fn main(@builtin(global_invocation_id) gid: vec3) { if (!uiResolvePixel(gid.xy)) { return; } @@ -452,10 +484,7 @@ fn main(@builtin(global_invocation_id) gid: vec3) { if (sp.x < lo.x || sp.y < lo.y || sp.x >= hi.x || sp.y >= hi.y) { continue; } let t = (sp - lo) / it.rect.zw; let uv = vec2(mix(it.uv.x, it.uv.z, t.x), mix(it.uv.y, it.uv.w, t.y)); - // stb_truetype SDF: pixel value ~128 is the edge. Treat alpha as - // the smoothed step around that midpoint. - let sdf = textureSampleLevel(atlasTex, atlasSampler, uv, 0.0).r; - let alpha = clamp((sdf - 0.5) * 8.0 + 0.5, 0.0, 1.0); + let alpha = uiGlyphCoverage(uv, it.uv, it.rect.zw); if (alpha <= 0.0) { continue; } let src = vec4(it.color.rgb, it.color.a * alpha); dst = uiBlendOver(dst, src); diff --git a/interfaces/Crafter.Graphics-FontAtlas.cppm b/interfaces/Crafter.Graphics-FontAtlas.cppm index 724df5d..e8ccd99 100644 --- a/interfaces/Crafter.Graphics-FontAtlas.cppm +++ b/interfaces/Crafter.Graphics-FontAtlas.cppm @@ -49,6 +49,11 @@ export namespace Crafter { static constexpr int kAtlasSize = 1024; static constexpr float kBaseSize = 32.0f; static constexpr int kPadding = 4; + // stb_truetype writes kOnEdgeValue + kPixelDistScale * distance (in + // atlas pixels) into each SDF byte. The text kernels decode it with + // copies of these two, UI_SDF_ON_EDGE and UI_SDF_PIXEL_DIST_SCALE in + // ui-shared.glsl and in the WGSL of dom-webgpu.js; the + // FontAtlasDecode test fails until all three agree. static constexpr int kOnEdgeValue = 128; static constexpr float kPixelDistScale = 32.0f; diff --git a/project.cpp b/project.cpp index 312f7b1..e6e32d2 100644 --- a/project.cpp +++ b/project.cpp @@ -420,6 +420,21 @@ extern "C" Configuration CrafterBuildProject(std::span a atlasFlushTest.config.files.emplace_back(fs::path("examples/HelloUI/font.ttf")); cfg.tests.push_back(std::move(atlasFlushTest)); + // Text from the UI kernels ramps over one screen pixel at any + // size and aspect. The kernels decoded FontAtlas's SDF with + // kPixelDistScale where 255 / kPixelDistScale belongs: a + // quarter-pixel ramp that made text look pixelated. Draws a + // string at 14 and 64 px and stretched, a glyph turned by its UVs + // and shrunk empty texels, through DispatchText and DispatchFused, + // and compares the frame with stb_truetype's own antialiased + // rasterizer. Loads the UI kernels and a font like + // UIFallbackParity. Needs the Wayland backend + a real compositor. + Test textAATest = LibraryTest("TextAntialiasing"); + textAATest.config.shaders = cfg.shaders; + textAATest.config.buildFiles = cfg.buildFiles; + textAATest.config.files.emplace_back(fs::path("examples/HelloUI/font.ttf")); + cfg.tests.push_back(std::move(textAATest)); + // Issue #160: a swapchain has as many images as its present mode // needs, and the frame loop takes any count. Every offered mode // is kept (no fallback to FIFO), and every frame presents what it @@ -628,6 +643,13 @@ extern "C" Configuration CrafterBuildProject(std::span a exampleShadersTest.requires_ = { "tool:glslang" }; cfg.tests.push_back(std::move(exampleShadersTest)); + // The text kernels decode FontAtlas's SDF with copies of its two + // encoding constants, in ui-shared.glsl and in the WGSL of + // dom-webgpu.js, which no test renders. Reads the sources and holds + // the copies to FontAtlas's, and checks that ui-text and ui-fused + // decode through the shared helper. No GPU. + cfg.tests.push_back(LibraryTest("FontAtlasDecode")); + // Issue #60: VulkanBuffer::FlushDevice / FlushHost now record the chosen // memory type's propertyFlags at Create time and skip the // flush/invalidate when the memory is HOST_COHERENT. The gate is pure diff --git a/shaders/ui-fused.comp.glsl b/shaders/ui-fused.comp.glsl index 0678dae..a6f5c81 100644 --- a/shaders/ui-fused.comp.glsl +++ b/shaders/ui-fused.comp.glsl @@ -46,11 +46,6 @@ layout(push_constant) uniform PC { layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in; -// SDF tuning for the text phase — must match Crafter::FontAtlas constants -// (same values as ui-text.comp.glsl). -const float ON_EDGE = 128.0 / 255.0; -const float DIST_SCALE = 32.0; - // Per-category clip-active bits packed into pc.misc.w (the fused analogue of // ui-shared.glsl's UI_FLAG_CLIP). A bit is set only when that category's clip // rect is narrower than the surface; when clear, the category skips its four @@ -79,9 +74,10 @@ shared uvec4 s_v4[UI_CHUNK]; // phases like the s_v* scratch (the per-pixel inner loops read them instead of // recomputing per pixel × item): // images: s_inv = 1.0/rect.zw -// text: s_inv = 1.0/rect.zw, s_band = SDF AA scale (a divide + two maxes) +// text: s_inv = 1.0/rect.zw, s_ramp = the SDF edge ramp (ui-shared.glsl's +// uiGlyphRamp, as in ui-text.comp.glsl) shared vec2 s_inv[UI_CHUNK]; -shared float s_band[UI_CHUNK]; +shared UIGlyphRamp s_ramp[UI_CHUNK]; shared uint s_keep[UI_CHUNK]; shared uint s_order[UI_CHUNK]; shared uint s_count; @@ -315,6 +311,10 @@ void main() { uint count = pc.itemCounts.w; bool inClip = inSurface && ((pc.misc.w & UI_FUSED_CLIP_TEXT) == 0u || uiPixelInClipRect(pxu, pc.clipText)); + // Queried only with text to draw: without a FontAtlas the font slots + // name nothing. + vec2 atlasSize = vec2(1.0); + if (count != 0u) atlasSize = uiGlyphAtlasSize(pc.misc.y, pc.misc.z); for (uint base = 0u; base < count; base += UI_CHUNK) { uint idx = base + lid; bool keep = false; @@ -323,12 +323,7 @@ void main() { s_v1[lid] = uiGlyphHeap[heap].items[idx].uv; s_v2[lid] = uiGlyphHeap[heap].items[idx].color; s_inv[lid] = 1.0 / s_v0[lid].zw; - // SDF AA band — atlas-px per screen-px, constant per glyph - // (uvSpan * kAtlasSize(1024) / screenSpan), max'd to 1px floor. - vec2 uvSpan = s_v1[lid].zw - s_v1[lid].xy; - vec2 atlasPerScreen = (uvSpan * 1024.0) * s_inv[lid]; - float scalePx = max(atlasPerScreen.x, atlasPerScreen.y); - s_band[lid] = max(scalePx, 0.0001); + s_ramp[lid] = uiGlyphRamp(s_v1[lid], s_inv[lid], atlasSize); keep = uiAabbOverlapsTile(s_v0[lid].xy, s_v0[lid].xy + s_v0[lid].zw, tileMin, tileMax); } @@ -349,18 +344,9 @@ void main() { vec2 t = (sp - s_v0[c].xy) * s_inv[c]; vec2 uv = mix(s_v1[c].xy, s_v1[c].zw, t); - // Font slots are push constants — provably dynamically uniform, - // so no nonuniformEXT (same as ui-text.comp.glsl). - float sdf = texture( - sampler2D(uiTextures[pc.misc.y], - uiSamplers[pc.misc.z]), - uv - ).r; - - float dAtlas = (ON_EDGE - sdf) * DIST_SCALE; - - // band (the AA scale) was precomputed once per glyph at load. - float a = clamp(0.5 - dAtlas / s_band[c], 0.0, 1.0); + // Font slots are push constants, dynamically uniform as + // uiGlyphCoverage needs (same as ui-text.comp.glsl). + float a = uiGlyphCoverage(pc.misc.y, pc.misc.z, uv, s_ramp[c], atlasSize); if (a <= 0.0) continue; vec4 col = s_v2[c]; diff --git a/shaders/ui-shared.glsl b/shaders/ui-shared.glsl index ace325d..88c48d1 100644 --- a/shaders/ui-shared.glsl +++ b/shaders/ui-shared.glsl @@ -157,6 +157,93 @@ float uiSdRoundRect(vec2 p, vec2 halfSize, vec4 r) { return min(max(q.x, q.y), 0.0) + length(max(q, 0.0)) - r.x; } +// ─── SDF text (FontAtlas) ─────────────────────────────────────────────── +// FontAtlas keeps each glyph as a signed distance field: stb_truetype writes +// kOnEdgeValue + kPixelDistScale * d into each texel, d being the distance +// from the outline in atlas pixels, positive inside, truncated to a byte and +// clamped to 0..255. These are FontAtlas's two constants in the encoder's own +// units; the FontAtlasDecode test holds them, and the WGSL copy in +// dom-webgpu.js, to FontAtlas's. ui-text and the fused kernel's text phase +// decode through the helpers below, and so should a copy of either. +const float UI_SDF_ON_EDGE = 128.0; // FontAtlas::kOnEdgeValue +const float UI_SDF_PIXEL_DIST_SCALE = 32.0; // FontAtlas::kPixelDistScale + +// The widest edge ramp, in atlas pixels, the stored distances can carry. A +// byte holds from (UI_SDF_ON_EDGE - 0.5) / UI_SDF_PIXEL_DIST_SCALE outside +// the outline to (255.5 - UI_SDF_ON_EDGE) / UI_SDF_PIXEL_DIST_SCALE inside, +// about 4 atlas pixels each way, and a ramp wider than twice the nearer end +// never reaches 0 there: below about 4 px every glyph quad's padding, all 0, +// would show as a translucent box. For text that small a ramp narrower than +// one screen pixel is the lesser evil. +const float UI_SDF_MAX_BAND = 2.0 * min(UI_SDF_ON_EDGE - 0.5, 255.5 - UI_SDF_ON_EDGE) + / UI_SDF_PIXEL_DIST_SCALE; + +// Distance from the outline in atlas pixels, negative inside, for an atlas +// sample (0..1). stb_truetype truncates, so byte b stands for [b, b + 1) of +// the encoded value and b + 0.5 is its unbiased middle: decoding b itself +// pulls every outline 1/64 atlas pixel inwards. +float uiSdfDistance(float sdf) { + return (UI_SDF_ON_EDGE - (sdf * 255.0 + 0.5)) / UI_SDF_PIXEL_DIST_SCALE; +} + +// What a glyph's per-pixel coverage needs, worked out once per glyph: +// atlasPerScreen the atlas pixels one screen pixel spans along x and y; +// coverage (k, b) of coverage = clamp(k * sdf + b, 0, 1) for a glyph +// scaled alike along x and y (within 1%), as ShapeText's +// always are: a ramp one screen pixel wide folded into one +// multiply-add. (0, 0) for one that isn't, whose ramp is +// then worked out per pixel. +struct UIGlyphRamp { + vec2 atlasPerScreen; + vec2 coverage; +}; + +// The atlas's size in texels: queried rather than taken from +// FontAtlas::kAtlasSize, so the two can't drift apart. +vec2 uiGlyphAtlasSize(uint texSlot, uint samplerSlot) { + return vec2(textureSize(sampler2D(uiTextures[texSlot], uiSamplers[samplerSlot]), 0)); +} + +// `uv` is the GlyphItem's (u0, v0, u1, v1), `invRectSize` 1 / its (w, h). +UIGlyphRamp uiGlyphRamp(vec4 uv, vec2 invRectSize, vec2 atlasSize) { + UIGlyphRamp ramp; + // abs: a glyph turned by swapping its UVs keeps its scale. + ramp.atlasPerScreen = abs(uv.zw - uv.xy) * atlasSize * invRectSize; + ramp.coverage = vec2(0.0); + float lo = min(ramp.atlasPerScreen.x, ramp.atlasPerScreen.y); + float hi = max(ramp.atlasPerScreen.x, ramp.atlasPerScreen.y); + if (hi <= lo * 1.01) { + // 0.5 - uiSdfDistance(sdf) / band, the band one screen pixel. + float inv = 1.0 / (UI_SDF_PIXEL_DIST_SCALE * clamp(hi, 0.0001, UI_SDF_MAX_BAND)); + ramp.coverage = vec2(255.0 * inv, 0.5 - (UI_SDF_ON_EDGE - 0.5) * inv); + } + return ramp; +} + +// The glyph's coverage at `uv` in the atlas. The slots must be dynamically +// uniform, as push constants are: there is no nonuniformEXT here, which lets +// the compiler hoist the descriptor load out of the per-pixel glyph loop. +float uiGlyphCoverage(uint texSlot, uint samplerSlot, vec2 uv, UIGlyphRamp ramp, vec2 atlasSize) { + float sdf = texture(sampler2D(uiTextures[texSlot], uiSamplers[samplerSlot]), uv).r; + if (ramp.coverage.x != 0.0) return clamp(ramp.coverage.x * sdf + ramp.coverage.y, 0.0, 1.0); + + // Scaled differently along x and y: one screen pixel spans + // atlasPerScreen.x atlas pixels across a vertical edge and + // atlasPerScreen.y across a horizontal one, so the ramp's width follows + // the direction the distance changes in, from one-texel differences. One + // width for both would be several pixels wide across one of them. + vec2 texel = 1.0 / atlasSize; + vec2 gradient = vec2( + texture(sampler2D(uiTextures[texSlot], uiSamplers[samplerSlot]), uv + vec2(texel.x, 0.0)).r, + texture(sampler2D(uiTextures[texSlot], uiSamplers[samplerSlot]), uv + vec2(0.0, texel.y)).r) - sdf; + float len = length(gradient); + // No gradient where the field is clamped, 4 atlas pixels from any + // outline, where the coverage is 0 or 1 at any width up to the cap. + float band = len > 0.0 ? length(gradient / len * ramp.atlasPerScreen) + : max(ramp.atlasPerScreen.x, ramp.atlasPerScreen.y); + return clamp(0.5 - uiSdfDistance(sdf) / clamp(band, 0.0001, UI_SDF_MAX_BAND), 0.0, 1.0); +} + // ─── shared-memory cooperative tile culling ───────────────────────────── // Each shader still dispatches one 8×8 workgroup per screen tile, but rather // than have all 64 threads walk the FULL item list — re-reading every item diff --git a/shaders/ui-text.comp.glsl b/shaders/ui-text.comp.glsl index 1b7afa9..3aa2dcf 100644 --- a/shaders/ui-text.comp.glsl +++ b/shaders/ui-text.comp.glsl @@ -21,20 +21,16 @@ layout(push_constant) uniform PC { layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in; -// SDF tuning — must match Crafter::FontAtlas::kOnEdgeValue / kPixelDistScale. -const float ON_EDGE = 128.0 / 255.0; -const float DIST_SCALE = 32.0; - shared vec4 s_rect[UI_CHUNK]; shared vec4 s_uv[UI_CHUNK]; shared vec4 s_color[UI_CHUNK]; // Per-glyph constants precomputed once at cooperative-load time and read by the // per-pixel inner loop, instead of being recomputed for every pixel × glyph: // s_invRectSize = 1.0/rect.zw (turns the per-pixel divide into a multiply) -// s_band = the SDF AA scale (a divide + two maxes that depend only on -// the glyph's uv span and rect size — fully loop-invariant) +// s_ramp = the SDF edge ramp (ui-shared.glsl's uiGlyphRamp), which +// depends only on the glyph's uv span and rect size shared vec2 s_invRectSize[UI_CHUNK]; -shared float s_band[UI_CHUNK]; +shared UIGlyphRamp s_ramp[UI_CHUNK]; shared uint s_keep[UI_CHUNK]; shared uint s_order[UI_CHUNK]; shared uint s_count; @@ -85,6 +81,10 @@ void main() { vec2 tileMin, tileMax; uiTileBounds(tileMin, tileMax); + // DispatchText draws only with a FontAtlas registered, so the font slots + // always name it here. + vec2 atlasSize = uiGlyphAtlasSize(pc.fontTextureSlot, pc.fontSamplerSlot); + uint lid = gl_LocalInvocationIndex; for (uint base = 0u; base < pc.hdr.itemCount; base += UI_CHUNK) { uint idx = base + lid; @@ -94,14 +94,7 @@ void main() { s_uv[lid] = uiGlyphHeap[pc.hdr.itemBuffer].items[idx].uv; s_color[lid] = uiGlyphHeap[pc.hdr.itemBuffer].items[idx].color; s_invRectSize[lid] = 1.0 / s_rect[lid].zw; - // Atlas-px per screen-px along this glyph's transform — keeps AA - // crisp at any rendering size. uvSpan * atlasSize / screenSpan, - // with FontAtlas::kAtlasSize = 1024. Constant for the whole glyph. - vec2 uvSpan = s_uv[lid].zw - s_uv[lid].xy; - vec2 atlasPerScreen = (uvSpan * 1024.0) * s_invRectSize[lid]; - float scalePx = max(atlasPerScreen.x, atlasPerScreen.y); - // 1-screen-px AA band, expressed in atlas-pixel units of dAtlas. - s_band[lid] = max(scalePx, 0.0001); + s_ramp[lid] = uiGlyphRamp(s_uv[lid], s_invRectSize[lid], atlasSize); keep = uiAabbOverlapsTile(s_rect[lid].xy, s_rect[lid].xy + s_rect[lid].zw, tileMin, tileMax); } @@ -123,20 +116,9 @@ void main() { vec2 t = (sp - s_rect[c].xy) * s_invRectSize[c]; vec2 uv = mix(s_uv[c].xy, s_uv[c].zw, t); - // Font slots are push constants — provably dynamically uniform, so no - // nonuniformEXT: lets the compiler hoist the descriptor load out of the - // per-pixel glyph loop and avoid the divergent-descriptor path. - float sdf = texture( - sampler2D(uiTextures[pc.fontTextureSlot], - uiSamplers[pc.fontSamplerSlot]), - uv - ).r; - - // Distance in atlas-pixels (negative inside the glyph). - float dAtlas = (ON_EDGE - sdf) * DIST_SCALE; - - // band (the AA scale) was precomputed once per glyph at load. - float a = clamp(0.5 - dAtlas / s_band[c], 0.0, 1.0); + // Font slots are push constants, dynamically uniform as + // uiGlyphCoverage needs. + float a = uiGlyphCoverage(pc.fontTextureSlot, pc.fontSamplerSlot, uv, s_ramp[c], atlasSize); if (a <= 0.0) continue; vec4 col = s_color[c]; diff --git a/tests/FontAtlasDecode/main.cpp b/tests/FontAtlasDecode/main.cpp new file mode 100644 index 0000000..e0c8ac9 --- /dev/null +++ b/tests/FontAtlasDecode/main.cpp @@ -0,0 +1,101 @@ +//SPDX-License-Identifier: LGPL-3.0-only +//SPDX-FileCopyrightText: Copyright (C) 2026 Catcrafts® + +// The text kernels decode FontAtlas's SDF with FontAtlas's own encoding. +// stb_truetype writes kOnEdgeValue + kPixelDistScale * distance into each +// byte, and the shaders carry copies of those two numbers: UI_SDF_ON_EDGE and +// UI_SDF_PIXEL_DIST_SCALE in shaders/ui-shared.glsl, which ui-text and the +// fused kernel decode through, and again in the WGSL text kernel of +// additional/dom-webgpu.js. The decode once multiplied by kPixelDistScale +// where 255 / kPixelDistScale belongs, which left every edge pixelated, and a +// copy that drifts from FontAtlas would go unseen on WebGPU, where no test +// renders (wasm runs none) and a small drift passes TextAntialiasing's +// tolerances anyway. So this reads the sources and compares the numbers, and +// checks that both GLSL kernels decode through the shared helper rather than +// arithmetic of their own. +// +// Pure CPU: reads files, no device. + +#include + +import Crafter.Graphics; +import std; + +using namespace Crafter; +namespace fs = std::filesystem; + +namespace { + +int failures = 0; + +void Check(bool ok, std::string_view what) { + std::println("{} {}", ok ? "PASS" : "FAIL", what); + if (!ok) ++failures; +} + +// The directory holding shaders/ and additional/, whether the test runs from +// its build output directory or the project root, as ExampleShaders finds it. +fs::path LocateRoot() { + for (const fs::path candidate : {".", "..", "../..", "../../.."}) { + std::error_code ec; + if (fs::is_regular_file(candidate / "shaders" / "ui-shared.glsl", ec) + && fs::is_regular_file(candidate / "additional" / "dom-webgpu.js", ec)) { + return candidate; + } + } + return {}; +} + +std::string ReadFile(const fs::path& path) { + std::ifstream in(path, std::ios::binary); + return std::string(std::istreambuf_iterator(in), std::istreambuf_iterator()); +} + +// The numbers `pattern`'s one capture group matches in `source`. +std::vector Matches(const std::string& source, const std::string& pattern) { + std::vector values; + const std::regex re(pattern); + for (auto it = std::sregex_iterator(source.begin(), source.end(), re); it != std::sregex_iterator(); ++it) { + values.push_back(std::stof((*it)[1].str())); + } + return values; +} + +// `declaration` is the source up to the name, `name` the constant, followed +// by "= ;" in either language. +void CheckConstant(std::string_view file, const std::string& source, const std::string& declaration, + std::string_view name, float expected) { + const std::vector values = Matches(source, declaration + R"(\s*=\s*([0-9]+(?:\.[0-9]+)?)\s*;)"); + std::string found = values.empty() ? "none" : ""; + for (std::size_t i = 0; i < values.size(); ++i) found += std::format("{}{}", i == 0 ? "" : ", ", values[i]); + Check(values.size() == 1 && values[0] == expected, + std::format("{} defines {} once, as {} (found {})", file, name, expected, found)); +} + +} // namespace + +int main() { + const fs::path root = LocateRoot(); + if (root.empty()) { + std::println(std::cerr, "could not find shaders/ and additional/ from {}", fs::current_path().string()); + return EXIT_FAILURE; + } + const std::string glsl = ReadFile(root / "shaders" / "ui-shared.glsl"); + const std::string wgsl = ReadFile(root / "additional" / "dom-webgpu.js"); + + const float onEdge = static_cast(FontAtlas::kOnEdgeValue); + const float pixelDistScale = FontAtlas::kPixelDistScale; + CheckConstant("ui-shared.glsl", glsl, R"(const\s+float\s+UI_SDF_ON_EDGE)", "UI_SDF_ON_EDGE", onEdge); + CheckConstant("ui-shared.glsl", glsl, R"(const\s+float\s+UI_SDF_PIXEL_DIST_SCALE)", "UI_SDF_PIXEL_DIST_SCALE", pixelDistScale); + CheckConstant("dom-webgpu.js", wgsl, R"(const\s+UI_SDF_ON_EDGE\s*:\s*f32)", "UI_SDF_ON_EDGE", onEdge); + CheckConstant("dom-webgpu.js", wgsl, R"(const\s+UI_SDF_PIXEL_DIST_SCALE\s*:\s*f32)", "UI_SDF_PIXEL_DIST_SCALE", pixelDistScale); + + for (const std::string_view kernel : {"ui-text.comp.glsl", "ui-fused.comp.glsl"}) { + const std::string source = ReadFile(root / "shaders" / kernel); + Check(source.find("uiGlyphCoverage(") != std::string::npos, + std::format("{} decodes the atlas through ui-shared.glsl's uiGlyphCoverage", kernel)); + } + + std::println("{}", failures == 0 ? "all checks passed" : "check(s) failed"); + return failures == 0 ? EXIT_SUCCESS : EXIT_FAILURE; +} diff --git a/tests/TextAntialiasing/main.cpp b/tests/TextAntialiasing/main.cpp new file mode 100644 index 0000000..0085873 --- /dev/null +++ b/tests/TextAntialiasing/main.cpp @@ -0,0 +1,486 @@ +//SPDX-License-Identifier: LGPL-3.0-only +//SPDX-FileCopyrightText: Copyright (C) 2026 Catcrafts® + +// Text drawn by the UI kernels has to be antialiased: a glyph's edge ramps +// from transparent to opaque over one screen pixel, at any size. FontAtlas +// keeps each glyph as an SDF in which stb_truetype writes kOnEdgeValue + +// kPixelDistScale * distance (in atlas pixels) per byte; the kernels decode a +// sample back into a distance and divide it by the atlas pixels one screen +// pixel spans along the glyph. ui-text and ui-fused used to multiply by +// kPixelDistScale instead of 255 / kPixelDistScale, which made the distance +// 1024/255 (about 4) times too large and the ramp a quarter of a pixel wide: +// nearly every edge pixel came out fully on or fully off, and text looked +// pixelated at every size. +// +// The test draws one scene through DispatchText and again through +// DispatchFused, captures the frame and checks: +// +// - one string at 14 and 64 px, below and above the atlas's 32 px base +// size, and at 14 px stretched to three times its width, each against +// stb_truetype's own antialiased rasterizer at the same size, aspect and +// sub-pixel position: about as many pixels are partly covered, and where +// the reference is partly covered the coverage is close to it. The +// quarter-pixel ramp gives a quarter as many; a ramp one atlas pixel +// wide, which ignores the glyph's size, half as many at 14 px and twice +// as many at 64 px; one ramp width for both axes of the stretched row, +// the width across its horizontal edges, makes it three pixels wide +// across the stems; +// - a glyph turned by swapping its UVs, the one way a GlyphItem flips, has +// to draw its mirror image. A signed UV span gave it hard edges; +// - atlas texels no glyph reaches hold 0, like every glyph's padding, and +// drawn shrunk far below 4 px they have to draw nothing. A ramp wider +// than the distances a byte can hold left the whole quad translucent; +// - the fused text phase draws what DispatchText draws. +// +// Needs a live Wayland compositor and a Vulkan device, like UIFallbackParity. + +#include "vulkan/vulkan.h" +#include "vulkan/vk_enum_string_helper.h" +#include + +// PNG decode for the captured frame. The library only compiles the writer +// (Window::CaptureNextFrame); STB_IMAGE_STATIC keeps the reader private to this file. +#define STB_IMAGE_IMPLEMENTATION +#define STB_IMAGE_STATIC +#define STBI_ONLY_PNG +#include "../../lib/stb_image.h" +// The reference rasterizer. Declarations only: the definitions are the +// library's own (Font.cpp), and Font::font is the stbtt_fontinfo they take. +#include "../../lib/stb_truetype.h" + +import Crafter.Graphics; +import Crafter.Event; +import std; + +using namespace Crafter; +namespace fs = std::filesystem; + +namespace { + +int failures = 0; + +void Check(bool ok, std::string_view what) { + std::println("{} {}", ok ? "PASS" : "FAIL", what); + if (!ok) ++failures; +} + +// The test binary's own directory: the build places the compiled shaders +// and font.ttf next to it, and the test runner's working directory is not +// that directory. +fs::path ExeDir() { + return fs::read_symlink("/proc/self/exe").parent_path(); +} + +constexpr std::uint32_t kWidth = 640; +constexpr std::uint32_t kHeight = 560; + +// Straight stems, curves and diagonals. +constexpr std::string_view kText = "Hamburgefonstiv"; +constexpr float kTextX = 20.0f; +constexpr std::uint32_t kGlyphCapacity = 64; + +struct Row { + float size; + float baseline; + float stretch; // horizontal scale applied to ShapeText's glyphs +}; +constexpr std::array kRows{{ {14.0f, 40.0f, 1.0f}, {64.0f, 120.0f, 1.0f}, {14.0f, 180.0f, 3.0f} }}; +// The glyph mirrored by its UVs, taken from the 64 px row, and where the +// original and the flipped copy go: whole pixels, so the check maps pixel to +// pixel. 'g' is neither left-right nor top-bottom symmetric. +constexpr char kMirrorGlyph = 'g'; +constexpr float kMirrorX = 340.0f; +constexpr float kMirrorY = 150.0f; +constexpr float kMirrorGap = 16.0f; +// The empty-texel item: 0.4 of the atlas drawn over 32 px is 12.8 atlas +// pixels per screen pixel, a glyph under 3 px. FontAtlas packs shelves from +// the top left, and this string fills one, so the bottom right is untouched. +constexpr float kEmptyX = 500.0f; +constexpr float kEmptyY = 150.0f; +constexpr float kEmptySize = 32.0f; +constexpr float kEmptyUv0 = 0.55f; +constexpr float kEmptyUv1 = 0.95f; +// The fused copy of the scene sits this much lower: a whole number of +// pixels, so it samples the atlas exactly where the standalone copy does. +constexpr int kFusedOffset = 270; + +// A pixel is an edge pixel when it is partly covered. +constexpr float kPartialLow = 0.02f; +constexpr float kPartialHigh = 0.98f; +// The SDF ramp is not the reference's box filter, so a correct kernel lands +// near the reference rather than on it: 1.02, 1.07 and 1.09 times its edge +// pixels (14 px, 64 px, 14 px stretched) on an RTX 4090, 1.02, 1.11 and 1.09 +// on llvmpipe, with a mean difference of 0.065 to 0.083 on them. The +// quarter-pixel ramp gives 0.31 and 0.21 times and 0.20; a ramp one atlas +// pixel wide 0.55 and 2.19 times and 0.14; one width for both axes 1.62 times +// on the stretched row. +constexpr float kMinEdgeRatio = 0.7f; +constexpr float kMaxEdgeRatio = 1.4f; +constexpr float kMaxEdgeError = 0.12f; + +template +void CreateItemBuffer(GraphicsBuffer& buffer, std::uint32_t count) { + buffer.Create( + VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, count); +} + +// Per-pixel coverage, 0 for the background and 1 for fully covered. +struct Coverage { + int width = 0; + int height = 0; + std::vector value; + + float At(int x, int y) const { return value[static_cast(y) * width + x]; } +}; + +int Code(float coverage) { + return static_cast(std::lround(coverage * 255.0f)); +} + +// White text over an opaque black background on an 8-bit UNORM swapchain: +// the red channel is the coverage the kernels composited. +Coverage LoadFrameCoverage(const fs::path& path) { + Coverage frame; + int channels = 0; + stbi_uc* pixels = stbi_load(path.string().c_str(), &frame.width, &frame.height, &channels, STBI_rgb_alpha); + if (pixels == nullptr) return frame; + frame.value.resize(static_cast(frame.width) * frame.height); + for (std::size_t i = 0; i < frame.value.size(); ++i) frame.value[i] = pixels[i * 4] / 255.0f; + stbi_image_free(pixels); + return frame; +} + +// stb_truetype's area coverage of kText laid out the way UIRenderer::ShapeText +// lays it out, stretched as the test stretches the row: the same pen positions +// (the atlas's advances, scaled and summed in the same order), each glyph +// rasterised at its own sub-pixel offset and composited "over" as the kernels +// do. +Coverage ReferenceCoverage(Font& font, const FontAtlas& atlas, const Row& row) { + Coverage ref{static_cast(kWidth), static_cast(kHeight), + std::vector(static_cast(kWidth) * kHeight, 0.0f)}; + const float scale = stbtt_ScaleForPixelHeight(&font.font, row.size); + const float atlasScale = row.size / FontAtlas::kBaseSize; + float cursor = 0.0f; + for (const char c : kText) { + const int codepoint = static_cast(c); + const Glyph* glyph = atlas.Lookup(font, static_cast(codepoint)); + if (glyph == nullptr) return ref; + const float pen = kTextX + cursor * row.stretch; + const float shift = pen - std::floor(pen); + const float scaleX = scale * row.stretch; + int x0 = 0, y0 = 0, x1 = 0, y1 = 0; + stbtt_GetCodepointBitmapBoxSubpixel(&font.font, codepoint, scaleX, scale, shift, 0.0f, &x0, &y0, &x1, &y1); + const int w = x1 - x0; + const int h = y1 - y0; + if (w > 0 && h > 0) { + std::vector bitmap(static_cast(w) * h); + stbtt_MakeCodepointBitmapSubpixel(&font.font, bitmap.data(), w, h, w, scaleX, scale, shift, 0.0f, codepoint); + const int left = static_cast(std::floor(pen)) + x0; + const int top = static_cast(row.baseline) + y0; + for (int y = 0; y < h; ++y) { + for (int x = 0; x < w; ++x) { + const int px = left + x; + const int py = top + y; + if (px < 0 || py < 0 || px >= ref.width || py >= ref.height) continue; + const float a = bitmap[static_cast(y) * w + x] / 255.0f; + float& dst = ref.value[static_cast(py) * ref.width + px]; + dst = a + dst * (1.0f - a); + } + } + } + cursor += glyph->advance * atlasScale; + } + return ref; +} + +// Half-open pixel box. +struct Box { + int x0 = 0, y0 = 0, x1 = 0, y1 = 0; +}; + +// The reference's covered pixels, plus a margin for the kernels' ramp, which +// reaches half a pixel past the outline. +Box CoveredBox(const Coverage& ref) { + Box box{ref.width, ref.height, 0, 0}; + for (int y = 0; y < ref.height; ++y) { + for (int x = 0; x < ref.width; ++x) { + if (ref.At(x, y) <= 0.0f) continue; + box.x0 = std::min(box.x0, x); + box.y0 = std::min(box.y0, y); + box.x1 = std::max(box.x1, x + 1); + box.y1 = std::max(box.y1, y + 1); + } + } + constexpr int kMargin = 2; + box.x0 = std::max(box.x0 - kMargin, 0); + box.y0 = std::max(box.y0 - kMargin, 0); + box.x1 = std::min(box.x1 + kMargin, ref.width); + box.y1 = std::min(box.y1 + kMargin, ref.height); + return box; +} + +bool Partial(float coverage) { + return coverage > kPartialLow && coverage < kPartialHigh; +} + +struct EdgeStats { + int frameEdges = 0; // partly covered pixels in the frame + int referenceEdges = 0; // partly covered pixels in the reference + float meanError = 0.0f; // mean |frame - reference| over the reference's edge pixels +}; + +// Compares the frame, `dy` pixels lower, with the reference inside `box`. +EdgeStats CompareEdges(const Coverage& frame, int dy, const Coverage& ref, const Box& box) { + EdgeStats stats; + double error = 0.0; + for (int y = box.y0; y < box.y1; ++y) { + for (int x = box.x0; x < box.x1; ++x) { + const float f = frame.At(x, y + dy); + const float r = ref.At(x, y); + if (Partial(f)) ++stats.frameEdges; + if (Partial(r)) { + ++stats.referenceEdges; + error += std::abs(f - r); + } + } + } + if (stats.referenceEdges > 0) stats.meanError = static_cast(error / stats.referenceEdges); + return stats; +} + +void CheckEdges(std::string_view what, const EdgeStats& stats) { + const float ratio = stats.referenceEdges > 0 + ? static_cast(stats.frameEdges) / static_cast(stats.referenceEdges) + : 0.0f; + Check(ratio >= kMinEdgeRatio && ratio <= kMaxEdgeRatio, + std::format("{}: {} antialiased edge pixels, stb_truetype's rasterizer {} ({:.2f} times, allowed {:.1f} to {:.1f})", + what, stats.frameEdges, stats.referenceEdges, ratio, kMinEdgeRatio, kMaxEdgeRatio)); + Check(stats.referenceEdges > 0 && stats.meanError <= kMaxEdgeError, + std::format("{}: coverage on the reference's edge pixels is {:.3f} from it on average (allowed {:.2f})", + what, stats.meanError, kMaxEdgeError)); +} + +std::string RowName(const Row& row) { + return row.stretch == 1.0f ? std::format("{} px", row.size) + : std::format("{} px stretched {}x", row.size, row.stretch); +} + +int Run() { + const fs::path dir = ExeDir(); + const fs::path outDir = fs::temp_directory_path() / "crafter-text-antialiasing"; + std::error_code ec; + fs::remove_all(outDir, ec); + fs::create_directories(outDir); + const fs::path png = outDir / "frame.png"; + + Device::Initialize(); + Window window(kWidth, kHeight, "TextAntialiasing test"); + // CaptureNextFrame writes 8-bit RGBA, swapping BGRA, and the coverage is + // read from those bytes. A surface whose 8-bit formats can't be storage + // images gets another format, which neither understands. + if (window.colorFormat != VK_FORMAT_B8G8R8A8_UNORM && window.colorFormat != VK_FORMAT_R8G8B8A8_UNORM) { + std::println("SKIP the swapchain format {} isn't 8-bit RGBA or BGRA, which the captured frame needs", + string_VkFormat(window.colorFormat)); + return 77; + } + VkCommandBuffer init = window.StartInit(); + + GraphicsDescriptorHeap heap; + heap.Initialize(/*images*/ 8, /*buffers*/ 8, /*samplers*/ 4); + window.descriptorHeap = &heap; + + Font font(dir / "font.ttf"); + FontAtlas atlas; + atlas.Initialize(init); + + UIRenderer ui; + ui.fontAtlas = &atlas; + ui.Initialize(window, heap, init, + dir / "ui-quads.comp.spv", dir / "ui-circles.comp.spv", + dir / "ui-images.comp.spv", dir / "ui-text.comp.spv", + dir / "ui-fused.comp.spv"); + window.passes.push_back(&ui); + + GraphicsBuffer background; + GraphicsBuffer standalone; + GraphicsBuffer fused; + CreateItemBuffer(background, 1); + CreateItemBuffer(standalone, kGlyphCapacity); + CreateItemBuffer(fused, kGlyphCapacity); + BufferSlot backgroundSlot = ui.RegisterBuffer(background); + BufferSlot standaloneSlot = ui.RegisterBuffer(standalone); + BufferSlot fusedSlot = ui.RegisterBuffer(fused); + + // Filled on the frame: the rows, then the mirrored pair and the + // empty-texel item. The fused buffer holds the same items kFusedOffset + // lower. + std::uint32_t glyphCount = 0; + std::array rowCounts{}; + GlyphItem original{}; + GlyphItem flipped{}; + EventListener buildSub(&ui.onBuild, [&](UIBuildArgs a) { + VkCommandBuffer cmd = a.cmd; + // The window never clears its image. Opaque black under white text + // makes the red channel the coverage. + background.value[0] = QuadItem{ 0, 0, 4096, 4096, 0.0f, 0.0f, 0.0f, 1.0f, 0, 0, 0, 0, 0, 0, 0, 0 }; + + GlyphItem* items = standalone.value; + glyphCount = 0; + std::uint32_t mirrorIndex = kGlyphCapacity; + for (std::size_t r = 0; r < kRows.size(); ++r) { + const Row& row = kRows[r]; + const std::uint32_t first = glyphCount; + rowCounts[r] = ui.ShapeText(font, row.size, kTextX, row.baseline, kText, {1.0f, 1.0f, 1.0f, 1.0f}, + items + first, kGlyphCapacity - first); + glyphCount += rowCounts[r]; + if (row.stretch != 1.0f) { + for (std::uint32_t i = first; i < glyphCount; ++i) { + items[i].x = kTextX + (items[i].x - kTextX) * row.stretch; + items[i].w *= row.stretch; + } + } + // Every character of kText has ink, so item i is character i. + if (row.size == 64.0f && row.stretch == 1.0f && rowCounts[r] == kText.size()) { + mirrorIndex = first + static_cast(kText.find(kMirrorGlyph)); + } + } + if (mirrorIndex < kGlyphCapacity && glyphCount + 3 <= kGlyphCapacity) { + original = items[mirrorIndex]; + original.x = kMirrorX; + original.y = kMirrorY; + flipped = original; + flipped.x = kMirrorX + original.w + kMirrorGap; + std::swap(flipped.u0, flipped.u1); + std::swap(flipped.v0, flipped.v1); + items[glyphCount++] = original; + items[glyphCount++] = flipped; + + GlyphItem empty{}; + empty.x = kEmptyX; + empty.y = kEmptyY; + empty.w = kEmptySize; + empty.h = kEmptySize; + empty.u0 = kEmptyUv0; + empty.v0 = kEmptyUv0; + empty.u1 = kEmptyUv1; + empty.v1 = kEmptyUv1; + empty.r = empty.g = empty.b = empty.a = 1.0f; + items[glyphCount++] = empty; + } + for (std::uint32_t i = 0; i < glyphCount; ++i) { + fused.value[i] = items[i]; + fused.value[i].y += kFusedOffset; + } + + constexpr VkAccessFlags read = VK_ACCESS_SHADER_READ_BIT; + constexpr VkPipelineStageFlags stage = VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT; + background.FlushDevice(cmd, read, stage); + standalone.FlushDevice(cmd, read, stage); + fused.FlushDevice(cmd, read, stage); + + ui.DispatchQuads(cmd, backgroundSlot, 1); + ui.DispatchText(cmd, standaloneSlot, glyphCount); + ui.DispatchFused(cmd, {}, {}, {}, {fusedSlot, glyphCount}); + }); + + // The first frame is enough: DispatchText and DispatchFused upload the + // glyphs ShapeText rasterised into the atlas before they draw, in the + // same command buffer, and the capture waits for the frame. + window.FinishInit(); + window.CaptureNextFrame(png); + window.Render(); + // The buffers, atlas and heap go out of scope while the frame may still + // be presenting; settle the queue first. + Device::CheckVkResult(vkQueueWaitIdle(Device::queue)); + + if (std::getenv("CRAFTER_GRAPHICS_VALIDATION") != nullptr) { + Check(Device::validationErrorCount == 0, + std::format("no Vulkan validation errors while rendering and capturing ({} seen)", Device::validationErrorCount)); + } + const bool shaped = std::ranges::all_of(rowCounts, [](std::uint32_t n) { return n == kText.size(); }); + Check(shaped && glyphCount == kText.size() * kRows.size() + 3, + std::format("every glyph was shaped ({} items for {} rows of {} glyphs and 3 more)", + glyphCount, kRows.size(), kText.size())); + + const Coverage frame = LoadFrameCoverage(png); + Check(!frame.value.empty(), "the frame was captured"); + if (frame.value.empty() || !shaped) return EXIT_FAILURE; + + // Both copies of the scene, one above the other, inside the frame. + const float sceneRight = std::max(kEmptyX + kEmptySize, flipped.x + flipped.w); + const float sceneBottom = std::max(kEmptyY + kEmptySize, flipped.y + flipped.h); + const bool fits = sceneRight <= static_cast(kWidth) && sceneBottom <= static_cast(kFusedOffset) + && frame.width >= static_cast(kWidth) && frame.height >= 2 * kFusedOffset; + Check(fits, std::format("the {}x{} frame holds both copies of the scene, {} by {} px each", + frame.width, frame.height, sceneRight, sceneBottom)); + if (!fits) return EXIT_FAILURE; + + for (const Row& row : kRows) { + const Coverage ref = ReferenceCoverage(font, atlas, row); + const Box box = CoveredBox(ref); + CheckEdges(std::format("DispatchText, {}", RowName(row)), CompareEdges(frame, 0, ref, box)); + CheckEdges(std::format("DispatchFused, {}", RowName(row)), CompareEdges(frame, kFusedOffset, ref, box)); + } + + // The flipped copy against the original read backwards on both axes. + // One code value of slack: the two sample the same atlas points through + // different float arithmetic. + { + const int w = static_cast(original.w); + const int h = static_cast(original.h); + const int ox = static_cast(original.x); + const int fx = static_cast(flipped.x); + const int y0 = static_cast(original.y); + int maxDiff = 0; + float peak = 0.0f; + for (int y = 0; y < h; ++y) { + for (int x = 0; x < w; ++x) { + const float a = frame.At(ox + x, y0 + y); + const float b = frame.At(fx + (w - 1 - x), y0 + (h - 1 - y)); + maxDiff = std::max(maxDiff, std::abs(Code(a) - Code(b))); + peak = std::max(peak, a); + } + } + Check(peak >= 0.99f && maxDiff <= 1, + std::format("a '{}' turned by swapping its UVs draws its mirror image (max difference {}, peak coverage {:.2f})", + kMirrorGlyph, maxDiff, peak)); + } + + // Texels holding 0 draw nothing, however far they are shrunk. + { + int maxCode = 0; + for (int y = static_cast(kEmptyY); y < static_cast(kEmptyY + kEmptySize); ++y) { + for (int x = static_cast(kEmptyX); x < static_cast(kEmptyX + kEmptySize); ++x) { + maxCode = std::max(maxCode, Code(frame.At(x, y))); + } + } + Check(maxCode == 0, std::format("empty atlas texels shrunk to {:.1f} atlas pixels per screen pixel draw nothing (max {})", + (kEmptyUv1 - kEmptyUv0) * FontAtlas::kAtlasSize / kEmptySize, maxCode)); + } + + // The whole fused copy against the whole standalone one. One code value + // of slack: the two kernels share the per-pixel math, but the driver + // compiles them separately and may contract or reorder it differently on + // an antialiased edge. + { + int maxDiff = 0; + for (int y = 0; y < kFusedOffset; ++y) { + for (int x = 0; x < static_cast(kWidth); ++x) { + maxDiff = std::max(maxDiff, std::abs(Code(frame.At(x, y)) - Code(frame.At(x, y + kFusedOffset)))); + } + } + Check(maxDiff <= 1, std::format("DispatchFused draws what DispatchText draws (max difference {})", maxDiff)); + } + + if (failures != 0) std::println(" frame kept for inspection: {}", png.string()); + return failures == 0 ? EXIT_SUCCESS : EXIT_FAILURE; +} + +} // namespace + +int main() { + const int result = Run(); + if (result != 77) std::println("{}", result == EXIT_SUCCESS ? "all checks passed" : "check(s) failed"); + return result; +}