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; +}