Watch

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 <noreply@anthropic.com>

2026-10-31 20:23:40 +01:00

20 KiB

Raw Blame History

Crafter.Graphics

Vulkan + WebGPU graphics library built around C++20 modules and bindless heaps. Provides window management, ray tracing, and a compute-shader-driven UI on a single, opinionated stack. Native builds use Vulkan with VK_EXT_descriptor_heap (the compute/UI path also runs without it, see below); wasm32-* builds target the browser via WebGPU and a DOM window backend.

Backends

Backends are chosen at build time by the target triple:

Target	Window	Renderer	Shaders
native Linux	Wayland	Vulkan (heap-bound)	GLSL → SPIR-V
native Windows	Win32	Vulkan (heap-bound)	GLSL → SPIR-V
`wasm32-*` (any)	DOM (canvas + JS env)	WebGPU	WGSL (loaded at runtime)

The two backends share the same C++ surface for the high-level pieces (UIRenderer, Mesh, RenderingElement3D, RTPass, item structs, FontAtlas, Image2D, ComputeShader). Backend-typed pieces (*Vulkan vs *WebGPU) live behind #ifdef CRAFTER_GRAPHICS_WINDOW_DOM. Vulkan ray tracing is hardware (VK_KHR_ray_tracing_pipeline); WebGPU ray tracing is a library-built software path (BVH + traceRay in a compute pipeline composed from user-supplied WGSL stages). The WebGPU path supports triangle and AABB (procedural, VK_GEOMETRY_TYPE_AABBS_KHR) geometry, closest-hit / miss / any-hit / intersection shaders — see examples/RTVolume for procedural spheres shaded through an intersection shader with an any-hit cut-out.

Native RT status: reading an acceleration structure through VK_EXT_descriptor_heap aborts with VK_ERROR_DEVICE_LOST on NVIDIA driver 610.43.02 — a driver-side fault in the brand-new descriptor-heap acceleration-structure path, not an engine bug (the setup is correct and validation-clean; images/buffers through the same heap work). The engine works around it transparently (issue #15): on the NVIDIA driver only, VulkanShader rewrites the compiled SPIR-V so heap AS reads become a TLAS-device-address + OpConvertUToAccelerationStructureKHR path (which reads no descriptor), and RTPass supplies the address as push data. Shaders and example code are unchanged, and it's a single fenced block gated on Device::workaroundDescriptorHeapAS, removable once a fixed driver ships. Still present on 615.78.08: with the rewrite off, traceRayEXT still loses the device and an inline ray query renders nothing. tests/HeapASRead covers the engine path; run HeapASRead --no-workaround to check a newer driver (it tries both paths with the rewrite off). Any app runs with the rewrite off while CRAFTER_GRAPHICS_NO_AS_WORKAROUND is set, to any value. See examples/VulkanTriangle/README.md for the full investigation. WebGPU RT is unaffected.

Native RT limitation (fixed in NVIDIA 615.78.08) — dynamic descriptor_heap indexing in hit shaders: on NVIDIA 610.43.02, indexing a descriptor_heap array with a runtime (non-constant) index inside a ray-tracing hit shader also device-losts (VK_ERROR_DEVICE_LOST), for plain SSBO and sampled-image descriptors. On 615.78.08 it works: tests/HeapDynamicIndexHit picks a buffer and a texture by a runtime index in a closest-hit, including different descriptors in neighbouring lanes of a wave. The rest of this note applies to the older drivers. A constant / spec-constant index is fine (that's why Sponza's closest-hit reads textures[albedoSlot] and samplers[samplerSlot] through spec constants), and the identical dynamic pattern works in fragment shaders (the UI renderer indexes uiTextures[] by per-item runtime slots) — so this is RT-stage-specific, not a general heap problem. Unlike the AS-read fault above this cannot be worked around transparently: sampled images have no device-address escape hatch the way an acceleration structure does (OpConvertUToAccelerationStructureKHR). The recommended pattern for bindless per-mesh geometry/material is to bind one resource and index within it dynamically rather than selecting a descriptor dynamically: pack geometry into a single SSBO (or reach it via buffer_reference) at a spec-constant slot and index by element offset, and put materials in one texture2DArray indexed by layer. Dynamic addressing inside a bound resource is ordinary memory/layer addressing and is unaffected; only dynamic selection of a descriptor faults. This is exactly what the WebGPU path already does (bucketed texture arrays + a single buffer). Full investigation and GLSL in examples/Sponza/README.md (issue #23). WebGPU RT is unaffected.

What's in here

Window — Wayland, Win32, and DOM backends, swapchain ring / canvas framing, input events. Pick a backend at build time via the target triple. The DOM backend routes every dynamic symbol through additional/dom-env.js and additional/dom-webgpu.js. On Wayland the window needs wl_compositor, wl_shm and xdg_wm_base, and uses xdg-decoration, wp_viewporter and wp_fractional_scale_v1 where the compositor has them. Without xdg-decoration (GNOME) the window has no title bar, as the library draws no client-side decorations; when either of the other two is missing it draws at scale 1, and the compositor scales it up on a HiDPI output. Natively the frame loop keeps three frames in flight (Window::numFrames) over however many images the swapchain has: as many as its present mode needs. Where the instance has surface_maintenance1 and the device swapchain_maintenance1, that is the mode's own count, which on Mesa's Wayland is below the surface's count for all modes (FIFO three, Mailbox four); without, the stricter of the two. window.currentFrame picks what lives per frame in flight (the command buffer, the descriptor-heap copy, RenderingElement3D::tlases, the frameIdx a pass is recorded with); window.currentImage is the swapchain image the frame draws into, window.images[window.currentImage]. A shader reads that image from a heap slot listed with heap.AddSwapchainSlot(slot, window), which the window rewrites with the acquired image every frame, so nothing is rewritten after a resize, and two windows can share a heap. A Resize or SetPresentMode from a listener of the frame being recorded takes effect after that frame's present.
Device (Vulkan only) — single-instance bring-up targeting VK_EXT_descriptor_heap; pipelines are created with VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT so there are no descriptor-set layouts and push constants travel via vkCmdPushDataEXT. Needs Vulkan 1.3. The heap, ray tracing, device-fault and swapchain-maintenance extensions are optional groups, detected at start-up (Device::descriptorHeapSupported, rayTracingSupported, deviceFaultSupported, swapchainMaintenance1Supported), as is the instance's surface_maintenance1 (surfaceMaintenance1Supported). Device features split by who uses them: the ones the binding model and the library's own shaders need are the library's to enable; arithmetic, storage and layout features an application's shaders use (shaderInt16 for a uint16_t spec constant, scalarBlockLayout, fp16, 16-bit storage) are the application's to ask for, as core feature structs: Device::Initialize({.vulkan10 = {.shaderInt16 = VK_TRUE}}). A second argument takes features to use where the GPU has them, and Device::enabledFeatures says which it got. A GPU that can't run the application is skipped. VulkanShader refuses a module that declares, directly or implicitly, a capability whose core feature, subgroup operation or optional group (ray tracing, the descriptor heap) is off, naming the field to pass or saying the GPU lacks it. Capabilities only an extension the library never enables could allow are left to the driver and the validation layer, as are layout rules (scalarBlockLayout), which have no capability.
DescriptorHeapVulkan / DescriptorHeapWebGPU — bindless slot allocators. Vulkan side allocates image/buffer/sampler slots in a VK_EXT_descriptor_heap; WebGPU side resolves slots to JS-side handle-table cookies that the dispatch bridge binds per pass. The window binds the Vulkan heap in every frame it records; a command buffer recorded outside the frame loop that dispatches against it calls heap.Bind(cmd, frame), which binds the frame's heaps, or its descriptor set on a driver without VK_EXT_descriptor_heap.
VulkanBuffer<T, Mapped> / WebGPUBuffer<T> — typed buffer. Vulkan variant has optional host mapping and a FlushDevice that issues the right host-write barrier; WebGPU variant goes through queue.writeBuffer over the JS bridge.
ImageVulkan<Pixel> / Image2D<Pixel> / Image2DArray<Pixel> — image + staging buffer with mip-chain support on Vulkan; on WebGPU, rgba8unorm 2D / 2D-array textures created and written via the bridge. Atlas (r8unorm, sub-region writes) is a separate path.
PipelineRTVulkan / PipelineRTWebGPU / ShaderBindingTableVulkan / ShaderBindingTableWebGPU / RTPass — ray-tracing pipelines. Vulkan uses native RT pipelines + SBTs; WebGPU compiles a wavefront / streaming software tracer — five @compute kernels (GENERATE → PREP → TRACE → SHADE → RESOLVE) sharing one module, connected by GPU ray/hit/payload buffers and a GPU-driven indirect bounce loop (dispatchWorkgroupsIndirect). TRACE carries zero user code (traversal + intersection only); user raygen calls rtEmitPrimaryRay, and closesthit / miss run in SHADE where they rtEmitRay continuation/shadow rays and rtAccumulate radiance (atomic — any number of rays per pixel per bounce, e.g. one shadow ray per light). An optional Resolve shader tonemaps the linear accumulator. See WAVEFRONT-DESIGN.md.
ComputeShader / WebGPUComputeShader — Tier 1 wrapper used by the UI system. Vulkan loads a .spv and dispatches with vkCmdPushDataEXT (or the .sets.spv build and vkCmdPushConstants without the heap); WebGPU loads a user-supplied .wgsl blob at runtime via wgpuLoadCustomShader. Use it directly for any custom compute.
UI — three-tier UI system; see below. The standard shaders (four kernels plus the fused one) ship as .spv / .sets.spv pairs on native and as WGSL strings baked into the WebGPU dispatcher.
FontAtlas — single-channel SDF atlas (1024×1024, 32pt base, shelf-packed, lazy Ensure per codepoint, dirty-flush via Update). Once it is full, a glyph that finds no free room draws blank for a frame but keeps its advance; when the next frame begins it gets a cell, rasterised then, taken from the glyphs drawn least recently but never one drawn in the last Device::framesInFlight + 1 frames (Device::frameCounter, which a Window advances). Evicting only then, and not while a frame is still asking for glyphs, keeps a frame from evicting glyphs it draws further down. A glyph taller than the shelves spans several. A GlyphItem is good for the frame it was shaped for: 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. The text kernels decode it through 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 per-vertex attribs) into global mesh heaps and builds the TLAS in a library compute pass.
Clipboard / Input / Gamepad / Router / Dom — input plumbing. Gamepad uses libudev+libevdev on Linux (libevdev is vendored in lib/libevdev) and WGI on Windows; the DOM backend exposes the host page DOM (Dom::HtmlElement) and a router for hash-routed wasm apps.

UI system (three tiers)

The UI is deliberately layered to balance no-boilerplate against no-lock-in:

Tier 1 — ComputeShader. Load any .spv, dispatch with push constants, library inserts inter-dispatch barriers. The escape hatch: if the standard shaders don't fit, write your own compute and dispatch it next to them.
Tier 2 — UIRenderer + standard shaders. Four shipped compute shaders (drawQuads, drawCircles, drawImages, drawText), POD item structs (QuadItem, CircleItem, ImageItem, GlyphItem), a shared GLSL contract in shaders/ui-shared.glsl, and helpers (RegisterBuffer, RegisterImage, RegisterSampler, FillHeader, Dispatch*, ShapeText). You build your own per-shader SSBOs (manual batching) and call one Dispatch* per shader type per frame. Item array order = draw order.
Tier 3 — stateless presentation functions. DrawButton, DrawCheckbox, DrawSlider, DrawProgressBar. Each is a small function that appends items to your buffers — they don't dispatch. Colors come in as small inline *Colors aggregates, no library Theme type. The source is the customization API: if a component doesn't fit, copy its body and edit it. No virtual hooks, no extension points.

What's not in the UI: widget tree, layout engine (just a Rect::SubRect carving helper), theming, hit-testing, focus management. State for interactive components (hover, drag, focus) lives in user-owned POD structs, not the library.

UI dispatch model

Standard shaders dispatch one workgroup per 8×8 screen tile — each thread iterates every item in the SSBO in array order, accumulating into a local dst, and stores once. Total cost is O(W·H·N); works well up to a few hundred items at 1080p. Splitting one buffer into multiple dispatches doesn't help — the same total work plus barrier overhead. If you need to render thousands of UI items, you want a different shader (tile binning, per-item-list resolve), not more dispatches.

Drivers without `VK_EXT_descriptor_heap`

The compute/UI path also runs on Vulkan drivers that lack VK_EXT_descriptor_heap (Mesa Turnip on Qualcomm Adreno, for one). Device::Initialize detects this and sets Device::descriptorHeapSupported = false. The same slots then live in one bindless descriptor set per frame instead of the heap, and push constants go through vkCmdPushConstants. Slot handles, item structs, UIDispatchHeader and the UIRenderer API are unchanged, so application code does not branch on it. Ray tracing is heap-only: on such a driver Device::rayTracingSupported is false and the RT entry points throw.

Each compute shader is shipped in two builds of the same source: <name>.spv for the heap and <name>.sets.spv for descriptor sets. ComputeShader::Load(path, heap) takes the .spv path and picks the right file itself. For a custom shader that means three things:

Declare bindless resources with the UI_HEAP_* macros from shaders/ui-shared.glsl, not with a literal descriptor_heap:
```
layout(UI_HEAP_BUFFERS, std430) readonly buffer MyBuf { MyItem items[]; } myHeap[];
```

cfg.shaders.emplace_back(fs::path("my.comp.glsl"), std::string("main"), ShaderType::Compute);
Shader& sets = cfg.shaders.emplace_back(fs::path("my.comp.glsl"), std::string("main"), ShaderType::Compute);
sets.defines = {"CRAFTER_GRAPHICS_DESCRIPTOR_SETS"};
sets.variant = "sets";

Keep the push-constant block within the device's maxPushConstantsSize (256 bytes on common drivers). Without the heap a larger payload throws in ComputeShader::Dispatch.

Set the environment variable CRAFTER_GRAPHICS_NO_DESCRIPTOR_HEAP to force the descriptor-set path on a driver that does have the heap; the UIFallbackParity test uses it to compare the two paths frame for frame.

Build

The repository is built with crafter-build (a project-config based build system; the project description lives in project.cpp):

crafter-build                        # native: Wayland on Linux, Win32 on Windows
crafter-build --target=wasm32-wasip1 # browser: DOM window + WebGPU renderer
crafter-build -r                     # build and run (in an example directory)

The build picks the window + renderer pair automatically from the target triple: any wasm32-* triple flips to DOM + WebGPU (no Vulkan loader linked), everything else stays on the native Vulkan path. Each example with both backends ships GLSL and WGSL copies of its shaders side-by-side (e.g. raygen.glsl + raygen.wgsl); project.cpp selects the right set per target.

Examples

See examples/. Quick map:

HelloWindow — minimal native window, no rendering.
HelloDom — wasm-only smoke test of the DOM partition: page-level events, HtmlElement::CreateInBody, and Router::PushState-driven SPA navigation. No GPU work.
VulkanTriangle — ray-traced triangle on both Vulkan and WebGPU. The smallest test of the bindless + RT path on each backend.
RTStress — wavefront RT benchmark: an N×N×N grid of a cube mesh (instance-count knob kGrid, 512 → 8000) shaded with primary + shadow rays. Prints a GPU timestamp-query per-pass breakdown each second. WebGPU/DOM only.
Sponza — ray-traced Sponza atrium on both backends. Exercises .cmesh / .ctex decompression (GPU VK_EXT_memory_decompression on Vulkan, CPU on WebGPU) and a textured closest-hit. See its README for asset provenance.
HelloUI — UI smoke test using all three tiers (background quad, slider, progress bar, button with text label, cursor-tracking circle).
CustomShader — Tier 1 demo: a user-authored compute shader inverting RGB under a list of item-circles, dispatched alongside the standard drawQuads. Shipped as both .comp.glsl and .comp.wgsl.
Decompression — Crafter::Compression CPU round-trip smoke test (used by the WebGPU asset path).
InputSystem — keyboard / mouse / gamepad event surface check.

License

This library is licensed under the GNU Lesser General Public License, version 3 only (LGPL-3.0-only). See LICENSE for the license text. The LGPL supplements the GNU General Public License version 3, a copy of which is included as LICENSE.GPL.

The example code under examples/ is licensed under the MIT license (see examples/LICENSE), so it can be freely copied into your own projects.

The lib/ directory contains vendored third-party code (stb, Wayland protocol bindings) that remains under its own permissive licenses; see the headers of those files.

20 KiB Raw Blame History Unescape Escape