Catcrafts/Crafter.Graphics
Watch
2
0
Fork
You've already forked Crafter.Graphics
 0
Code Issues 25 Pull requests Projects Releases Packages Wiki Activity Actions

docs: AGENTS.md for working on the library, AGENTS.consumer.md for its consumers

Browse source 
AGENTS.md
- Guidance for agents and contributors working on Crafter.Graphics: the
  layout and the partitions by backend, the dev environment and the
  targets, the build and test commands, how tests are declared and what
  each group needs, code style, the design rules the code embodies,
  commit conventions, security notes and gotchas. It carries the
  dependency-notes section crafter-build 3.0.0 writes, holding
  Crafter.Math's notes.

AGENTS.consumer.md
- What an agent working in a consumer needs: depending on the library,
  the program shape and the frame order, device features, lifetimes, the
  heap and its descriptor-set fallback, buffers and images, the UI tiers,
  custom compute shaders, ray tracing, input, the DOM build, the
  environment variables, testing a consumer, and which example to copy.
  crafter-build injects it into consumers' AGENTS.md on their builds.

Both were reviewed against the code. The review found that --timing no
longer compiles (#178), which both files note.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
Jorijn van der Graaf 2026-10-31 16:02:25 +01:00
commit 576977bf35
Signed by: jorijnvdgraaf
GPG key ID: 2937E59CDCC1BCFB
2 changed files with 370 additions and 0 deletions
Download patch file Download diff file Expand all files Collapse all files

164
AGENTS.consumer.md Normal file
View file

@ -0,0 +1,164 @@
## Crafter.Graphics
The Crafter engine's graphics library: a C++ module, `import Crafter.Graphics;`, that opens a window, drives a Vulkan 1.3+ device through `VK_EXT_descriptor_heap` (bindless: no descriptor-set layouts, push data through `vkCmdPushDataEXT`), traces rays through `VK_KHR_ray_tracing_pipeline`, draws UI from compute shaders, and reads keyboards, mice, touch, gamepads and the clipboard. A `wasm32-*` build of the same surface runs in the browser on WebGPU with a software ray tracer. Everything is in `namespace Crafter`. The vectors and matrices are Crafter.Math's (`import Crafter.Math;`) and the compressed asset formats Crafter.Asset's (`import Crafter.Asset;`); both ship notes of their own. Window callbacks are `Crafter::Event<T>` from Crafter.Event (`import Crafter.Event;`): `EventListener<void> sub(&window.onMouseMove, [&] { ... });` subscribes, and the listener's destructor unsubscribes. Source, examples and issues: https://forgejo.catcrafts.net/Catcrafts/Crafter.Graphics (LGPL-3.0-only; the examples under `examples/` are MIT and meant to be copied). These notes are its `AGENTS.consumer.md`: a fix to them goes there, not into the copy crafter-build writes.
### Building against it
- The target picks the backend, in the library and in your code. Any `wasm32-*` target builds the DOM + WebGPU library; everything else builds Vulkan, with Wayland on Linux (there is no X11 backend) or Win32 on `x86_64-w64-mingw32`. The library defines `CRAFTER_GRAPHICS_WINDOW_WAYLAND`, `_WIN32` or `_DOM` for its own units only, so a wasm consumer sets `cfg.target = "wasm32-wasip1"`, `cfg.defines.push_back({"CRAFTER_GRAPHICS_WINDOW_DOM", ""})`, `cfg.compileFlags.push_back("-msimd128")` (the dependency PCMs were compiled with it) and `EnableWasiBrowserRuntime(cfg)`, as `examples/HelloUI/project.cpp` does. Portable code `#ifdef CRAFTER_GRAPHICS_WINDOW_DOM`s at the few sites where the backends differ; the library has no shared signatures that are no-ops on one side.
- The library's link flags (the Vulkan loader, `wayland-client`, `xkbcommon`, `udev`, or the Win32 libraries) come with the dependency. Its shaders (`ui-*.comp.spv` and `ui-*.comp.sets.spv`) and, for wasm, `dom-env.js` and `dom-webgpu.js` are copied into your executable's bin directory, as are your `cfg.shaders` outputs (`<name>.spv` from `<name>.glsl`) and `cfg.files`. Paths are resolved against the working directory at runtime, so run from the bin directory, as `crafter-build -r` does.
- Your own shaders: `cfg.shaders.emplace_back(fs::path("raygen.glsl"), std::string("main"), ShaderType::RayGen)`, with `Compute`, `RayGen`, `Miss`, `ClosestHit`, `AnyHit`, `Intersect` and `Callable`. A compute shader for the UI heap is registered twice, see "Custom compute shaders". WGSL for the DOM build ships through `cfg.files` and is loaded at runtime.
- Project args the library understands, passed through your `args`: `--local` (sibling checkouts of Crafter.Event, Crafter.Math and Crafter.Asset instead of clones), `--timing` (meant to add frame timing counters; the build does not compile with it at the moment, issue #178), `--no-webgpu` (wasm: a DOM-only page without the WebGPU bridge, which otherwise requests a GPU adapter on every load and replaces the page body with an error where WebGPU is missing).
- Needs at runtime: Vulkan 1.3 (1.4 is used where the GPU has it), a presentable surface, and the features the library's own shaders use. `VK_EXT_descriptor_heap` and ray tracing are optional groups, see "Device and features". Mesa's llvmpipe runs the compute/UI path, and recent Mesa advertises the heap.
### A program, and what the loop does
```cpp
Device::Initialize({.vulkan10 = {.shaderInt16 = VK_TRUE}}); // the features your shaders use; on the DOM build Initialize() takes none
Window window(1280, 720, "Title"); // `static` on the DOM build
VkCommandBuffer init = window.StartInit(); // a one-time setup command buffer
DescriptorHeapVulkan heap;
heap.Initialize(/*images*/ 16, /*buffers*/ 16, /*samplers*/ 4);
window.descriptorHeap = &heap; // bound at the start of every frame
// create buffers, images, meshes, pipelines and passes, recording into `init` ...
window.passes.push_back(&pass); // RenderPass*: a UIRenderer, an RTPass, your own
window.FinishInit(); // submits `init` and waits for it
window.Render(); // one frame
window.StartUpdate(); // frames on the compositor's clock from here
window.StartSync(); // the event loop; never returns
```
- Include `vulkan/vulkan.h` yourself, before the imports (the examples wrap it in `#ifndef CRAFTER_GRAPHICS_WINDOW_DOM`): the library includes it only in its own global module fragments, and neither Vulkan's types nor its macros (`VK_TRUE`, the `VK_*` flags) cross an `import`. On the DOM build `Device::Initialize()` takes no arguments and `DeviceFeatures` does not exist, which is why the dual-target examples keep a separate DOM `main`.
- `StartSync` runs the platform loop until `open` is false (`onClose` fires first). Each pass of the loop dispatches platform events into the `Window` events, runs `Gamepad::Tick()` and `onBeforeUpdate`, and runs `Update()` for a window that is due. On Wayland a window on the compositor's clock updates from its frame callback during that event dispatch, before the pass's `onBeforeUpdate`, and the loop wakes for every event, so `onBeforeUpdate` fires several times per frame; on Win32 and in the browser `onBeforeUpdate` precedes the same pass's `Update`. `updateInterval`, set before `StartUpdate`, makes a window update every so often whether visible or not; pair it with `SetPresentMode(PresentMode::Mailbox)` so a present never waits for the display. `AddPollFd(fd, callback)` (Wayland) wakes the loop for a socket or an eventfd.
- `Update()` computes `mouseDelta` and calls `Render()`, which acquires a swapchain image, waits for that image's fence (three frames in flight, `Window::numFrames`), reclaims deferred deletions, begins `drawCmdBuffers[currentBuffer]`, binds the heap (or the descriptor set), fires `onUpdate(FrameTime)`, records every pass in `passes` in order with a barrier on the swapchain image between them, transitions the image to present, submits and presents. Per-frame GPU work of your own (a TLAS rebuild, a physics dispatch) is recorded from `onUpdate` into `window.drawCmdBuffers[window.currentBuffer]`; the heaps are bound by then and the passes come after. `onResize` fires after a swapchain rebuild; anything holding descriptors of `window.imageViews[]` rewrites them there.
- A `RenderPass` is `Record(cmd, frameIdx, window)` plus `SwapchainStage()`, which must return the stage that touches the swapchain image (`COMPUTE_SHADER`, `RAY_TRACING_SHADER` or `TRANSFER`): the frame's barriers take their stage and access masks from it. The image stays in `VK_IMAGE_LAYOUT_GENERAL` through the whole pass list, with `STORAGE | TRANSFER_SRC` usage and `TRANSFER_DST` where the surface allows it (`window.swapchainUsage`). Nothing clears it: draw a background quad, or clear from a transfer pass where `TRANSFER_DST` is there.
- `window.width` and `height` are physical pixels, as is `currentMousePos`. `CaptureNextFrame(path)` saves the next frame as a PNG. `SetTitle`, `SetAppId`, `SetCursorImage`, `SetDefaultCursor`; `SetTextInput(active, x, y, w, h, purpose)` tells the compositor's input method a text field has focus, which opens an on-screen keyboard on a phone.
- Several windows can exist (on Wayland `Device::windows` lists them; the other backends keep no list), but they share one frame clock: one rendering window is the assumption behind the deletion queue and the staging rings.
### Device and features
- `Device::Initialize(required = {}, optional = {})` creates the instance, weighs every GPU and creates the device. `DeviceFeatures` holds the five core feature structs, `vulkan10` to `vulkan14`; only their `VkBool32`s are read, so `{.vulkan12 = {.shaderFloat16 = VK_TRUE}}` is a complete argument. The library enables what its own shaders use and nothing more: a `uint16_t` spec constant needs `.vulkan10.shaderInt16`, fp16 arithmetic `.vulkan12.shaderFloat16`, 16-bit storage `.vulkan11.storageBuffer16BitAccess`, scalar layout `.vulkan12.scalarBlockLayout`. A GPU lacking a required feature is skipped; when none is left the call throws naming each GPU's gaps, and may be called again asking for less. Optional features are enabled where the GPU has them; `Device::enabledFeatures` says which, `supportedFeatures` what the GPU has, `apiVersion` the version in use (`vulkan14` features need a 1.4 GPU; `maintenance5` is always on). `CRAFTER_GRAPHICS_DEVICE=<n>` picks GPU `n`.
- `VulkanShader` checks every module's SPIR-V capabilities against what was enabled and throws naming the file and the field to pass (`vulkan10.shaderInt16`), or saying the GPU lacks it. Most drivers accept such a module silently, so this is the error you get here instead of a wrong frame on another machine.
- Optional groups, decided once at `Initialize`: `Device::descriptorHeapSupported` (`VK_EXT_descriptor_heap`), `rayTracingSupported` (needs the heap too), `memoryDecompressionSupported` (`VK_EXT_memory_decompression`: `.cmesh` and `.ctex` decompress on the GPU; otherwise the library decodes them on the CPU), `deviceFaultSupported`. The RT entry points throw through `Device::RequireRayTracing` without ray tracing; the compute/UI path runs everywhere.
- `Device::device`, `queue`, `commandPool`, `physDevice`, `instance`, `memoryProperties`, `deviceProperties`, `rayTracingProperties` and `descriptorHeapProperties` are public statics for raw Vulkan. `Device::GetMemoryType(typeBits, required, preferred = 0)` and `PreferDirectDeviceWrite(bytes)` are the allocation helpers.
- `Device::pipelineCache` is shared by every pipeline and written to `Device::pipelineCachePath` (`pipeline_cache.bin` in the working directory) at exit. Set the path before `Initialize`, and gitignore the file.
### Lifetimes
Frames are pipelined three deep and nothing waits for the GPU at steady state, so:
- Every library object that owns a Vulkan handle (`VulkanBuffer`, `Mesh`, `TlasWithBuffer`, `VulkanShader`, `ComputeShader`, `PipelineRTVulkan`, and `ImageVulkan`'s staging buffer) is move-only and releases the handle through `Device::EnqueueDeletion`; the render loop frees it `Device::framesInFlight` frames later. Destroying one mid-frame is safe, and so is destroying it from another thread. `VulkanBuffer::Clear()` frees at once and is only for a GPU known to be idle; `DeferredClear()` is the queued form. A program with no `Window` reclaims nothing on its own: call `Device::DrainDeletions()` after `vkQueueWaitIdle(Device::queue)`.
- Static owners (a `static Mesh`, a `static VulkanBuffer`) are fine: their destructors run after the driver has unloaded, and the queue allows for that by never being destroyed itself. `RenderingElement3D::tlases` is leaked for the same reason.
- `Window` is neither movable nor copyable, and its destructor waits for the device, so a window can be closed and another opened.
- `ImageVulkan` has no destructor: call `Destroy()` once the GPU is done with it. `FontAtlas` waits for the queue and destroys its own image.
- Heap slots (`ImageSlot`, `BufferSlot`, `SamplerSlot`) are RAII handles that return the slot on destruction, and a returned slot is handed out again at once. Free a slot only when no frame in flight reads it (after a `vkQueueWaitIdle`, or at a scene switch), and before its heap goes. In descriptor-set mode a slot write destroys the previous view immediately: same rule.
- `RenderingElement3D::elements` holds raw pointers: `Add`/`Remove` elements whose storage does not move (reserve the vector, or make them `static`). `UIRenderer::ShapeText` and `FontAtlas` cache by `Font*`: call `ui.InvalidateFont(font)` before destroying a `Font`.
### The bindless heap
`DescriptorHeapVulkan::Initialize(images, buffers, samplers)` sizes one resource heap and one sampler heap per frame (`resourceHeap[Window::numFrames]`, `samplerHeap[...]`, mapped `VulkanBuffer`s) and the slot allocators. `AllocateImageSlots(n)`, `AllocateBufferSlots(n)` and `AllocateSamplerSlots(n)` return `{firstElement, count}` ranges (contiguous; freed runs are reused), `Free*Slots` returns them, and running out throws naming the pool and the counts.
- **Buffer slot indices are relative.** `AllocateBufferSlots` indexes the heap's buffer region, which starts after the image region, while a GLSL `descriptor_heap` SSBO array indexes buffer descriptors from heap byte 0. The shader-visible index is `heap.bufferStartElement + raw`. The RAII `BufferSlot` converts to exactly that (its `operator std::uint16_t()`), `UIRenderer::RegisterBuffer` returns one, and raw code passes `heap.bufferStartElement` as a spec constant: the RT examples declare `layout(constant_id = 0) const uint16_t bufferStart = 0us;` and index `topLevelAS[bufferStart]`, which is why they ask for `shaderInt16`. Image slots index the heap as they are, in image-descriptor units; samplers live in their own heap. A wrong buffer index reads zeros, which looks like every field of every item being zero. In descriptor-set mode `bufferStartElement` is 0, so code that goes through the conversion is right in both.
- Writing a descriptor by hand: `Device::vkWriteResourceDescriptorsEXT` into `heap.resourceHeap[f].value + heap.ImageByteOffset(imageRange.firstElement)` or `heap.BufferByteOffset(bufferRange.firstElement)`, `Device::descriptorHeapProperties.imageDescriptorSize` or `bufferDescriptorSize` bytes long, for each frame `f`, then `heap.resourceHeap[f].FlushDevice()`. `BufferByteOffset` takes the raw index `AllocateBufferSlots` returned and adds the region start itself; handing it a `BufferSlot`, which converts to the absolute index, adds the start twice and puts the descriptor in the wrong place. `ImageByteOffset` takes an image slot as it is. Per-frame resources (the swapchain image `window.imageViews[f]`, `tlases[f]`) are written into frame `f`'s heap; a shared resource into every frame's. `examples/VulkanTriangle/main.cpp` is the pattern. The `UIRenderer::Register*` helpers do this for you in both binding models.
- **Descriptor-set fallback.** On a driver without `VK_EXT_descriptor_heap` (Mesa Turnip on Adreno) `Device::descriptorHeapSupported` is false and the same slots live in one bindless descriptor set per frame; slot handles, item structs, `UIDispatchHeader` and the `UIRenderer` API are unchanged, so application code does not branch. Ray tracing is unavailable there. `CRAFTER_GRAPHICS_NO_DESCRIPTOR_HEAP` forces this path on a driver that has the heap, to test it. Raw `vkWriteResourceDescriptorsEXT` code is heap-only; a consumer that must run on such a driver stays on the `UIRenderer::Register*` helpers.
### Buffers, images, samplers
- `VulkanBuffer<T, Mapped>`: `Create(usage, memoryFlags, count, preferredFlags = 0)`; `Resize(...)` reuses the allocation when the request fits and the usage matches, otherwise queues the old one and creates; `size` (bytes), `buffer`, `address` (non-zero only with `VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT`), `memoryPropertyFlagsChosen`. With `Mapped`, `value` is the `T*` of a persistent map: after writing, call `FlushDevice()` (whole buffer), `FlushDevice(offset, bytes)` (a range, rounded to `nonCoherentAtomSize`) or `FlushDevice(cmd, dstAccess, dstStage)`, which flushes and records the host-write barrier; all three are no-ops on coherent memory. `FlushHost()` is the invalidate for readback. A UI item buffer is `GraphicsBuffer<QuadItem, true>` with `STORAGE_BUFFER | SHADER_DEVICE_ADDRESS` usage and `HOST_VISIBLE` memory with `DEVICE_LOCAL` as the preferred fourth argument (the examples pass both as required, which throws on a GPU without a host-visible device-local type), flushed each frame before the dispatch that reads it.
- Static geometry: `VulkanBuffer<T, false>::UploadDeviceLocal(usage, src, count, cmd, dstAccess, dstStage)` places the buffer in device memory, written directly on ReBAR and UMA and staged through a per-frame staging ring on a small BAR, and records the barrier to `dstStage`. `UploadDeviceLocalRange(src, offset, count, cmd, dstAccess, dstStage)` re-uploads the dirty window `[offset, offset + count)`; `src` points at the first dirty element (`vertices.data() + offset`), not at the array's start. A re-upload into a buffer a frame may still read is ordered after those reads on the GPU.
- `ImageVulkan<Pixel>` (`Image2D<Pixel>` is the portable alias): `Create(width, height, mipLevels, cmd, format, usage, layout, consumerStage = RAY_TRACING_SHADER, streamed = false)` makes the image and a mapped staging buffer `image.buffer`; fill `image.buffer.value` and `Update(cmd, layout)` copies it in, builds the mip chain by blits and transitions to `layout` for `consumerStage`. `Update(const CompressedTextureAsset&, cmd, layout)` decompresses a `.ctex` on the GPU, or the CPU. A static image releases its staging after the first upload; `streamed = true` keeps it (the font atlas, through `UpdateRegion`). Pass `VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT` as `consumerStage` for an image a compute shader samples; the RT default is corrected to compute on a device without ray tracing. `image`, `imageView` and `descriptor` are public for raw descriptor writes; `Destroy()` releases.
- `SamplerVulkan<Pixel>` makes one linear, repeat, 16x anisotropic sampler at construction (`textureSampler`). The UI has `UIRenderer::RegisterSampler(const VkSamplerCreateInfo&)` and `RegisterLinearClampSampler()`.
### UI
Three tiers, on purpose: Tier 1 is any compute shader you write (`ComputeShader`, `UIRenderer::Dispatch`); Tier 2 the four standard kernels, their item structs and the `UIRenderer` helpers; Tier 3 stateless functions that append items to your buffers. The library has no widget tree, layout engine, theme, hit-testing or focus management: state lives in your PODs, hit-testing is `Rect::Contains` against the rects you drew, and a Tier 3 component that does not fit is copied out of `implementations/Crafter.Graphics-UIComponents.cpp` and edited, not subclassed.
Setup: `FontAtlas atlas; atlas.Initialize(init);` then `UIRenderer ui; ui.fontAtlas = &atlas; ui.Initialize(window, heap, init); window.passes.push_back(&ui);`. `Initialize` loads the five kernels by file name from the working directory (`ui-quads.comp.spv` and so on; the `.sets.spv` builds must sit beside them) and registers the swapchain images and the atlas. Then, per item buffer, `auto slot = ui.RegisterBuffer(buf);` gives a `BufferSlot` ready for a shader.
Per frame, in an `EventListener<UIBuildArgs>` on `ui.onBuild` (fired from the pass's `Record` with `{cmd, frameIdx}`):
1. Rebuild the item arrays in your mapped buffers, in draw order: later items overdraw earlier ones. Lay out with `Rect::FromWindow(window)`, `SubRect(size, Anchor)` and `Inset(...)`.
2. `buf.FlushDevice(cmd, VK_ACCESS_SHADER_READ_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT)` for each buffer with items.
3. Dispatch: `ui.DispatchQuads(cmd, slot, count, clipRectPx = {0, 0, 1e9, 1e9})`, `DispatchCircles`, `DispatchImages`, `DispatchText`, or `ui.DispatchFused(cmd, {quadsSlot, qc}, {circlesSlot, cc}, {imagesSlot, ic}, {glyphsSlot, gc})`, which composites the four categories in that fixed order (quads, circles, images, text) with one load and one store; an empty category costs nothing. The two forms mix; a custom `ui.Dispatch` between two `DispatchFused` calls is the flush point.
- 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.
- 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<T>(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.
### Custom compute shaders (Tier 1)
A GLSL compute shader that dispatches beside the standard kernels follows `shaders/ui-shared.glsl`, which the build puts on the include path:
```glsl
#version 460
#extension GL_GOOGLE_include_directive : enable
#include "ui-shared.glsl"
struct MyItem { vec4 a; }; // mirrored in C++ with a static_assert on its size
layout(UI_HEAP_BUFFERS, std430) readonly buffer MyBuf { MyItem items[]; } myHeap[];
layout(push_constant) uniform PC { UIDispatchHeader hdr; /* your fields after it */ } pc;
layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in;
void main() {
ivec2 px;
if (!uiResolveScreenPixel(pc.hdr, px)) return;
// read myHeap[pc.hdr.itemBuffer].items[i]; imageLoad/imageStore uiImages[pc.hdr.outImage]
}
```
- 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 `<name>.spv` and `<name>.sets.spv`. `ComputeShader::Load("<name>.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.
- 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.
### Ray tracing (Vulkan)
Heap-only: check `Device::rayTracingSupported`, or expect `std::runtime_error` from the first RT call. `examples/VulkanTriangle/main.cpp` is the whole setup.
- Shaders: `VulkanShader shader("raygen.spv", "main", VK_SHADER_STAGE_RAYGEN_BIT_KHR, &specInfo);` (`nullptr` for no specialisation); `ShaderBindingTableVulkan table; table.Init(shaders);`; `PipelineRTVulkan pipeline; pipeline.Init(cmd, raygenGroups, missGroups, hitGroups, table, maxRecursionDepth = 1);` with spans of `VkRayTracingShaderGroupCreateInfoKHR` indexing the shader array. **`maxRecursionDepth` counts the raygen as 1**: a closest-hit that traces a shadow ray needs 2, and tracing past the pipeline's depth loses the device.
- Geometry: `Mesh mesh; mesh.Build(std::span<Vector<float, 3, 3>> vertices, std::span<std::uint32_t> indices, cmd, RTBuildOptions{});` builds a BLAS and sets `mesh.blasAddr`; `Build(const CompressedMeshAsset&, cmd)` builds from a `.cmesh`, decompressed on the GPU where supported; `BuildProcedural(aabbs, opaque, cmd)` builds from `RTAabb` boxes for a `VK_RAY_TRACING_SHADER_GROUP_TYPE_PROCEDURAL_HIT_GROUP_KHR` with an intersection shader, and has an overload taking a device address a compute pass wrote. `RTBuildOptions` is `.preference` (`FastTrace`, `FastBuild`), `.allowUpdate` and `.noDuplicateAnyHit`; with `allowUpdate`, `Refit(vertices, indices, cmd)` and its dirty-range overload update in place and keep `blasAddr`. A rebuild into storage that still fits keeps `blasAddr` too; read it again after a rebuild that grows.
- Instances: a `RenderingElement3D` holds a `VkAccelerationStructureInstanceKHR instance` (`transform`, `mask`, `instanceCustomIndex`, `instanceShaderBindingTableRecordOffset`, `flags`, `accelerationStructureReference = mesh.blasAddr`; Crafter.Math's `MatrixRowMajor<float, 4, 3, 1>::Store(reinterpret_cast<float*>(instance.transform.matrix))` writes the transform), `userMetadata` (copied per instance id into `tlases[f].metadataBuffer`) and `transformOwnedByGpu`. `RenderingElement3D::Add(&e)` and `Remove(&e)` register it. `RenderingElement3D::BuildTLAS(cmd, frameIndex, preference = FastTrace)` builds or refits `tlases[frameIndex]` from every registered element: once per frame slot at setup (`0`, `1`, `2`), and per frame from `onUpdate` for `window.currentBuffer` when instances move. `e.MarkHostDataDirty()` after a change to an element's host fields lets the per-frame copy skip unchanged slots. `tlases[f].address` goes into a `VK_DESCRIPTOR_TYPE_ACCELERATION_STRUCTURE_KHR` descriptor of frame `f` (see "The bindless heap") and changes only when the TLAS outgrows its storage.
- The pass: `RTPass rtPass(&pipeline); rtPass.pushPtr = &camera; rtPass.pushBytes = sizeof(camera);` pushes your data at offset 0 of the shaders' push block before every trace (the bytes are read at record time, so updating `camera` before the frame is enough); `width` and `height` trace at a size other than the window's. `window.passes.push_back(&rtPass);`.
- In GLSL, with the extensions the examples enable (`GL_EXT_ray_tracing`, `GL_EXT_descriptor_heap`, `GL_EXT_nonuniform_qualifier` for a non-constant heap index, `GL_EXT_shader_explicit_arithmetic_types_int16` for a `uint16_t` spec constant, and `GL_EXT_shader_image_load_formatted` or a `writeonly` image): `layout(descriptor_heap) uniform accelerationStructureEXT topLevelAS[];`, `layout(descriptor_heap) uniform writeonly image2D image[];`, indices from spec constants or push data. The output image is the swapchain image at its image slot, in `VK_IMAGE_LAYOUT_GENERAL`. Start from `examples/VulkanTriangle/raygen.glsl`.
- **NVIDIA.** On the proprietary driver (still on `615.78.08`) reading an acceleration structure through the heap loses the device. The library works around it transparently: `VulkanShader` rewrites the SPIR-V so the heap AS read becomes a TLAS-address load from the push block plus `OpConvertUToAccelerationStructureKHR`, and `RTPass` and `ComputeShader` push the address. Nothing changes in your shaders or code, except that the rewrite appends an 8-byte member after your push block, and that it turns `shaderInt64` on for itself only: a shader of yours that declares `Int64` still has to ask for `shaderInt64`, or it runs on NVIDIA alone. `CRAFTER_GRAPHICS_NO_AS_WORKAROUND` disables the rewrite, to probe a new driver. Second note: on `610.43.02`, indexing a `descriptor_heap` array with a runtime value inside a hit shader faulted (fixed in `615.78.08`). The portable pattern, which the Sponza example uses, is one descriptor at a spec-constant slot and dynamic indexing inside it: all geometry in one SSBO indexed by element offset, materials in one `texture2DArray` indexed by layer. `examples/Sponza/README.md` has the GLSL.
### Input
- `Window` events, `Event<void>` unless noted: `onMouseMove` (then `currentMousePos`, `lastMousePos`, `mouseDelta`), `onMouseEnter` and `onMouseLeave` with `mouseInside` (native only: the DOM build declares the two events but never fires them, and has no `mouseInside`), `onMouseLeftClick`, `onMouseLeftHold`, `onMouseLeftRelease` and the right-button trio (`mouseLeftHeld`, `mouseRightHeld`), `onMouseScroll` (`Event<std::uint32_t>`: whole detents, reinterpret the payload as `int32`, positive is wheel down, every backend normalised to one per notch), `onRawKeyDown`, `onRawKeyHold`, `onRawKeyUp` (`Event<KeyCode>`, raw platform codes; `heldKeys`), `onTextInput` (`Event<const std::string_view>`, UTF-8; repeats are synthesised on Wayland), `onPreedit` and `onTextInputDelete` (input methods), `onTouchDown`, `onTouchMove`, `onTouchUp` (`Event<TouchPoint>`; the first finger also drives the pointer events), `onClose`, `onResize`, `onBeforeUpdate`, `onUpdate` (`Event<FrameTime>`: `now`, `delta`).
- Keys: `Key(CrafterKeys::Space)` is `consteval` and folds to the platform's raw code (a Win32 scancode, a Linux keycode), so bindings stay portable in source and compare against `onRawKeyDown`'s `KeyCode` at runtime. Bindings serialised with `Input::BindingToString` are platform-specific.
- `Input::Map`: `auto& jump = map.AddAction("Jump", Input::ActionType::Button);` (`Button`, `Axis1D`, `Vector2`), `jump.bindings = { Input::KeyBind{Key(CrafterKeys::Space)}, Input::GamepadButtonBind{id, Gamepad::Button::South} };` (also `MouseButtonBind`, `MouseScrollBind`, `MouseDeltaBind{scale}`, `GamepadAxisBind`, `GamepadStickBind{id, Gamepad::Stick::Left}`, `WASDBind{up, down, left, right}`), `map.Attach(window);`, and `map.Tick();` once per frame, after `Gamepad::Tick()`, which the loop runs before `onBeforeUpdate`. Every `Tick` recomputes the mouse delta and drains the scroll accumulator, and on Wayland `onBeforeUpdate` fires on every loop wake, so either tick from `onUpdate`, or tick in `onBeforeUpdate` and consume the delta and scroll actions in that same listener, as the Sponza example does. Read `action.pressed`, `value` or `vector2`, or subscribe to `onPerformed`, `onCanceled`, `onValueChanged`, `onVector2Changed`. `StartRebind(action, CaptureMask, callback)` captures the next input.
- `Gamepad::connected`, `onConnected`, `onDisconnected`, per-device `buttons[]`, `axes[]` and events, `Rumble(dev, low, high, duration)`, `FindById(id)`; buttons are named by position (`South`, `East`, `West`, `North`). `Clipboard::SetText(text)` and `GetText()` (`std::optional<std::string>`).
### WebGPU and the DOM build
The same C++ surface, with these differences:
- The `Window` is the page, and `Device::Initialize()` takes no arguments. Make the window `static`: `StartSync` hands control to `requestAnimationFrame` and `_Exit(0)`s without running static destructors, so nothing after it runs on any target. `StayAlive` is `StartSync` without the animation loop, for pages that only touch the DOM, and is `StartSync` natively. `Window::numFrames` is 1, so `tlases[0]` and `BuildTLAS(cmd, 0)`. `GraphicsCommandBuffer` is an opaque `std::uint32_t`.
- `GraphicsBuffer<T, true>::Create(count)`, `FlushDevice()` without a command buffer (`queue.writeBuffer`), `FlushDeviceRange`, `FlushDeviceSlot`, and `EnqueueReadback` with `PollReadback` for GPU to CPU. `Image2D<T>::Create(w, h)` is `rgba8unorm`, `Update(const CompressedTextureAsset&)` decompresses on the CPU, `AllocateSlot(heap)` registers it; `Image2DArray<T>` stacks layers with a CPU mip chain; `StorageImage2D` is a GPU-written float texture. `AllocateLinearClampSampler(heap)` and `AllocateLinearRepeatSampler(heap)`.
- `DispatchImages` has a six-argument overload with one texture and one sampler for the whole dispatch (WebGPU bind groups); the four-argument overload binds the font atlas instead. `DispatchFused` issues the four kernels in order through the four-argument overloads, so quads, circles and text match the Vulkan result but an images category does not: draw images with the six-argument `DispatchImages` on this backend.
- Custom compute in the UI pass: `WebGPUComputeShader::Load(wgslPath, bindings, rayQuery = false)`, where each `UICustomBinding` is `{group >= 2, binding, kind, format, pushOffset}` with `kind` one of `Buffer`, `SampledTexture`, `Sampler`, `SampledTextureArray`, `BufferReadWrite`, `StorageTexture`; the renderer reads the slot at `pushOffset` out of your push data and binds that handle. Group 0 binding 0 is the header uniform and group 1 the ping-pong images, both declared by your WGSL; with `rayQuery` true the library injects its own group 1 (the RT heap and the `rayQuery*` API) and your shader must not declare one. Standalone compute: `PlainComputeShader::Load(wgsl, pushUniformSize, bindings, rayQuery)` and `Dispatch(push, bytes, handles, gx, gy, gz)` run anywhere in the frame, `handles` parallel to `bindings`. It has no library bind groups: group 0 binding 0 is your optional push uniform (`pushUniformSize` bytes, 0 for none) and your bindings start at group 1, or at group 2 with `rayQuery`, which reserves group 1 the same way.
- Ray tracing is a library-built wavefront tracer composed from your WGSL stages: `WebGPUShader(fs::path("raygen.wgsl"), entryFn, WebGPURTStage::Raygen | Miss | ClosestHit | AnyHit | Intersection | Resolve)` (wrap the path: a bare string literal is ambiguous with the constructor that takes the source itself), with the entry signatures documented in `interfaces/Crafter.Graphics-ShaderBindingTableWebGPU.cppm`; `ShaderBindingTableWebGPU::Init`, `RTShaderGroup{.type = RTShaderGroupType::General | TrianglesHitGroup | ProceduralHitGroup, ...}`, `PipelineRTWebGPU::Init(cmd, raygen, miss, hit, sbt, bindings, hdrOutputFormat)`. Raygen calls `rtEmitPrimaryRay`; closest-hit and miss run in the SHADE stage and call `rtEmitRay` (continuation and shadow rays) and `rtAccumulate` (atomic, any number of rays per pixel). `RTPass` adds `maxDepth` (bounces), `raysPerPixel` (the ray buffer budget), `handlesPtr` and `handlesCount` (your bindings' handles) and `outTexHandle` (HDR output). `Mesh::Build`, `BuildProcedural` and `RenderingElement3D` keep their API, but `Refit` and the host-span `RefitProcedural` rebuild the host BVH under a new `blasAddr`: set `instance.accelerationStructureReference` again before the next `BuildTLAS`, or portable deforming-mesh code keeps tracing the old geometry here. Only the device-buffer `RefitProcedural` keeps the handle.
- `Crafter::Dom::HtmlElement` and `HtmlElementPtr` (move-only DOM handles with typed listeners), `Crafter::Router` (history, location, query parsing) and `Crafter::Dom::*Event` exist only here. `WebGPU::SetCanvasMount(elementId)` renders into an element instead of the page.
- Examples with a DOM build ship `.glsl` and `.wgsl` side by side; `RTStress`, `RayQueryPick`, `HDRBloom` and `RTMultiShadow` are WebGPU-only.
### Environment variables
All are "set to any value", read once at `Device::Initialize` or swapchain creation.
| Variable | Effect |
|---|---|
| `CRAFTER_GRAPHICS_VALIDATION` | The Khronos validation layer with GPU-assisted validation, plus info messages; off by default. `Device::validationErrorCount` counts the errors. |
| `CRAFTER_GRAPHICS_DEVICE=<n>` | Use GPU number `n` as Vulkan lists them, never another; an error when it can't run the application. |
| `CRAFTER_GRAPHICS_NO_DESCRIPTOR_HEAP` | Take the descriptor-set path on a driver that has the heap (and lose ray tracing), as a heap-less driver would. |
| `CRAFTER_GRAPHICS_NO_AS_WORKAROUND` | Turn the NVIDIA SPIR-V rewrite off, to check a new driver. |
| `CRAFTER_GRAPHICS_NO_VULKAN_1_4` | Create a Vulkan 1.3 instance on a 1.4 GPU. |
| `CRAFTER_PRESENT_IMMEDIATE` | Present immediately, uncapped, for benchmarks. |
### Testing a consumer
`Device::Initialize` on Linux connects to Wayland right after creating the instance, so even a windowless test needs a compositor (`WAYLAND_DISPLAY`). `Window::CaptureNextFrame(png)` followed by `Render()` writes a frame to read back and compare; decode it with a private copy of `stb_image.h` (`STB_IMAGE_STATIC`), since the library compiles only the writer. Assert `Device::validationErrorCount == 0` only when `CRAFTER_GRAPHICS_VALIDATION` is set, as the count stays zero without the layer. Return 77 (crafter-build's skip) when `!Device::rayTracingSupported` and the test needs it. A test with GPU work and no `Window` calls `vkQueueWaitIdle(Device::queue)` and `Device::DrainDeletions()` before it ends. The library's own `tests/` has one of each kind to copy.
### Examples to copy from
`HelloWindow` (a window), `HelloUI` (all three UI tiers, text, hit-testing), `CustomShader` (a Tier 1 shader in GLSL and WGSL, registered twice), `VulkanTriangle` (the smallest RT setup, raw descriptor writes), `Sponza` (`.cmesh` and `.ctex` through `Mesh::Build` and `Image2D`, texture arrays, a free camera on `Input::Map`), `RTVolume` (procedural geometry with intersection and any-hit shaders), `InputSystem` (actions, rebinding, gamepads), `HelloDom` (DOM and router, no GPU), `Decompression` (the asset round trip). All MIT.

206
AGENTS.md Normal file
View file

@ -0,0 +1,206 @@
# AGENTS.md
Guidance for coding agents and contributors working on Crafter.Graphics itself. The [README](README.md) is the user-facing manual, the example READMEs hold the driver investigations, and [AGENTS.consumer.md](AGENTS.consumer.md) is what agents working in a *consumer* of this library read: crafter-build copies it into the dependency-notes section of the consumer's `AGENTS.md` on every build. When any of them disagrees with the code, the code wins; the interfaces in `interfaces/` are the reference.
## Project overview
Crafter.Graphics is the Crafter engine's graphics library: one C++ module, `Crafter.Graphics`, with 41 partitions, built with [crafter-build](https://forgejo.catcrafts.net/Catcrafts/Crafter.Build) from `project.cpp`. On native targets it is a Vulkan 1.3+ renderer built around `VK_EXT_descriptor_heap` (bindless, no descriptor-set layouts, push data through `vkCmdPushDataEXT`), with a bindless-descriptor-set fallback for drivers without the heap, hardware ray tracing, a compute-shader UI, a Wayland or Win32 window, and input. On `wasm32-*` targets the same C++ surface runs in a browser over WebGPU through a JS bridge, with a software wavefront ray tracer. It depends on Crafter.Event (events), Crafter.Math (vectors, matrices) and Crafter.Asset (`.cmesh`, `.ctex`), all pulled from `forgejo.catcrafts.net` unpinned, and, on native targets, on the Khronos `Vulkan-Headers` and `Vulkan-Utility-Libraries` from GitHub as header-only `ExternalDependency` entries. Its consumers are the examples under `examples/` (each a `LocalProject` of `../../project.cpp`) and the maintainer's applications, which pull `master`.
| Path | What lives there |
|---|---|
| `interfaces/Crafter.Graphics.cppm` | The umbrella: `export import`s every partition, unconditionally. |
| `interfaces/Crafter.Graphics-<Partition>.cppm` | One partition per file, see below. The Vulkan header and the platform headers (`wayland-client.h`, `windows.h`, `stb_truetype.h`, the Linux input headers) go in the global module fragment behind the backend's `#ifdef`. |
| `implementations/Crafter.Graphics-<Partition>.cpp` | Implementation units, in two sets chosen by `project.cpp`: 15 native and 17 DOM units. `Clipboard`, `Font`, `FontAtlas`, `Gamepad`, `Input`, `InputField`, `UI-Shared`, `UIComponents` and `Window` are in both; `-Device-Features.cpp` is a second unit of `:Device`; `-UI-WebGPU.cpp`, `-Mesh-WebGPU.cpp` and `-RenderingElement3D-WebGPU.cpp` are the DOM halves; `-GamepadWin32.c` and `.h` are the C binding to Windows.Gaming.Input. |
| `shaders/` | The five standard UI compute kernels (`ui-quads`, `ui-circles`, `ui-images`, `ui-text`, `ui-fused`) and `ui-shared.glsl`, the GLSL contract user shaders `#include`. Each kernel is compiled twice, see "Design rules". |
| `additional/dom-env.js`, `additional/dom-webgpu.js` | The JS half of the DOM backend: the wasm `env` imports for the window, DOM, router and clipboard, and the WebGPU bridge with the WGSL of the UI kernels, the TLAS build and the wavefront tracer baked in. Shipped as `cfg.files` of every wasm build. |
| `lib/` | Vendored: stb_image, stb_image_write, stb_truetype; the generated Wayland protocol C files (xdg-shell, xdg-decoration, fractional-scale, viewporter, text-input-v3); `libevdev/` with its kernel headers. Not edited. |
| `examples/<Name>/` | Standalone projects, MIT, each with its own `project.cpp`, `main.cpp` and shaders in GLSL and WGSL side by side. `examples/README.md` is the index. |
| `tests/<Name>/main.cpp` | One test per directory, plus the GLSL and the `font.ttf` some need. Declared in `project.cpp`. |
| `project.cpp` | The build description: the dependencies, the backend switch on the target, the partition and implementation lists, the shader builds, libevdev, and every test. |
| `README.md`, `examples/*/README.md` | The manual, and the NVIDIA driver investigations (`VulkanTriangle`, `Sponza`). |
| `AGENTS.consumer.md` | Notes for consumers' agents; crafter-build injects them into consumers' `AGENTS.md`. |
| `bin/`, `build/`, `pipeline_cache.bin` | Generated and gitignored. `build/test-logs/` keeps the output of failed tests; `pipeline_cache.bin` is written by any example or test run from here. |
Partitions, by what they are for. Every partition file exists on every target; one that does not apply compiles to an empty module behind `#ifdef CRAFTER_GRAPHICS_WINDOW_DOM` (or `#ifndef`), and several hold both halves in one file (`Window`, `Device`, `UI`, `Mesh`, `RenderingElement3D`, `RTPass`, `RenderPass`, `Image2D`, `FontAtlas`, `RT`).
- Both backends: `:Types` (`FrameTime`, `TouchPoint`, `CrafterKeys`, `AlignUp`), `:Keys` (`consteval Key(CrafterKeys)` to the platform's raw code), `:Window`, `:Device`, `:Input` (named actions, rebinding), `:Gamepad`, `:Clipboard`, `:Animation`, `:Font` (stb_truetype), `:FontAtlas` (the SDF atlas), `:UI` (`UIRenderer`, the item structs, `Rect`), `:UIComponents` (Tier 3), `:InputField`, `:RenderPass`, `:RTPass`, `:Mesh`, `:RenderingElement3D`, `:Image2D`, `:RT` (portable RT PODs and flag constants), `:GraphicsTypes` (the `using` aliases `GraphicsBuffer`, `GraphicsDescriptorHeap`, `GraphicsComputeShader`, `GraphicsCommandBuffer`), `:ForwardDeclarations`.
- Vulkan only: `:VulkanBuffer`, `:ImageVulkan`, `:SamplerVulkan`, `:DescriptorHeapVulkan` (heap and descriptor-set mode, the RAII slots), `:ComputeShader` (Tier 1), `:ShaderVulkan` (`VulkanShader` and the NVIDIA SPIR-V rewrite), `:ShaderBindingTableVulkan`, `:PipelineRTVulkan`, `:VulkanTransition`, `:Decompress` (`VK_EXT_memory_decompression`).
- DOM only: `:WebGPU` (the `extern "C"` bridge declarations), `:WebGPUBuffer`, `:DescriptorHeapWebGPU`, `:WebGPUComputeShader`, `:PlainComputeShader`, `:ShaderBindingTableWebGPU`, `:PipelineRTWebGPU`, `:Dom`, `:DomEvents`, `:Router`.
## Dev environment
crafter-build is the only build tool: the Flatpak (`alias crafter-build='flatpak run net.catcrafts.CrafterBuild'`; a locally built one runs with `--branch=local`) or a native build of the sibling checkout. `crafter-build --version` prints the compiler it resolved. Loading `project.cpp` clones Crafter.Event, Crafter.Math and Crafter.Asset into the crafter-build cache and, since they are unpinned, `git pull`s them on every invocation, and clones the two Khronos header repositories, so a build needs the network. `--local` takes the sibling checkouts (`../Crafter.Event`, `../Crafter.Math`, `../Crafter.Asset`) instead and hands the flag on, so Crafter.Asset picks the sibling Crafter.Math too; two copies of one module on the path do not link.
The host is Linux on Wayland. The native Linux backend is Wayland only (no X11), so every test that calls `Device::Initialize` needs `WAYLAND_DISPLAY`: it connects to the display right after creating the Vulkan instance, window or no window. The build links `wayland-client`, `xkbcommon`, `udev` and the Vulkan loader from the host or the sysroot; libevdev is vendored and built as its own static library, `Crafter.Graphics-libevdev`, once per build. A Vulkan 1.3 driver is the floor; `VK_EXT_descriptor_heap` (with `shaderUntypedPointers`) and ray tracing are optional groups detected at start-up. Recent Mesa's llvmpipe advertises the heap, so a machine without a GPU runs the UI path. The Khronos validation layer is optional and off by default.
Targets, picked with `--target=` and passed down to the dependencies:
| Target | Backend | Notes |
|---|---|---|
| native (`x86_64-unknown-linux-gnu`) | Wayland + Vulkan | the development loop |
| `x86_64-w64-mingw32` | Win32 + Vulkan | cross-compiled; links `vulkan-1`, `kernel32`, `user32`, `gdi32`, `runtimeobject`; gamepads through Windows.Gaming.Input |
| `wasm32-wasip1` | DOM + WebGPU | no Vulkan loader, no Wayland; `dom-env.js` and `dom-webgpu.js` ship beside the `.wasm`; `--no-webgpu` drops the WebGPU bridge for DOM-only apps |
| `aarch64-alpine-linux-musl` with `--sysroot=`, `--march=armv8-a --mtune=generic` | Wayland + Vulkan | the Mesa Turnip phone target: no descriptor heap, no ray tracing, the descriptor-set path |
Project flags, before or after `test`: `--local`, `--timing` (defines `CRAFTER_TIMING` for `Window::LogTiming` and per-listener update timings; the gated code references `Window` members that no longer exist, so the build fails, issue #178), `--no-webgpu` (wasm only), and crafter-build's `--target=`, `--march=`, `--mtune=`, `--sysroot=`, `--debug`. Flags `ApplyStandardArgs` does not interpret are part of the variant hash, so each combination has its own `bin/Crafter.Graphics-<target>-<march>-<mtune>-<hash>/` with `libCrafter.Graphics.a`, the PCMs and the shaders (`ui-*.comp.spv` and `ui-*.comp.sets.spv`); earlier variants stay until `clean`. Glob for a variant directory rather than predicting its hash.
The Flatpak sandbox sees the host filesystem but not the host `PATH`: the host's binaries are under `/run/host/usr/bin`, from where statically linked tools (`wine`, `qemu-aarch64-static`) run and dynamically linked ones often do not, and it has its own `/tmp`. A `--runner=cmd:<tool>` is probed with `which` inside the sandbox; a missing runner fails those tests rather than skipping them, and skips their build.
## Build, test and style commands
```bash
crafter-build # the native library and its shaders
crafter-build --local # against ../Crafter.Event, ../Crafter.Math, ../Crafter.Asset
crafter-build --target=wasm32-wasip1 # the DOM/WebGPU library
crafter-build --target=x86_64-w64-mingw32 # the Win32 library
crafter-build test # build and run every test for the host
crafter-build test 'TLAS*' # name glob
crafter-build test --list
crafter-build test --jobs=1 --timeout=120
CRAFTER_GRAPHICS_VALIDATION=1 crafter-build test # with the validation layer; the tests' validation checks only count then
cd examples/HelloUI && crafter-build -r # build an example against this checkout and run it from its bin dir
crafter-build clean
```
`project.cpp` registers no lint rules, so `crafter-build lint` and `format` have nothing to do here; style is kept by hand (see "Code style"). There is no CI. Before a change is done: `crafter-build test` passes on the host with `CRAFTER_GRAPHICS_VALIDATION` set; the examples whose API you touched build (`cd examples/<Name> && crafter-build`), since they are the in-tree consumers; `crafter-build --target=wasm32-wasip1` still compiles when you touched a partition with a DOM half or a backend-agnostic one; and `crafter-build --target=x86_64-w64-mingw32` still compiles when you touched `Window`, `Device`, `Gamepad` or `Clipboard`.
## Testing instructions
Tests are declared in `project.cpp`, not discovered, through the `LibraryTest(name)` lambda: it copies the library's settings field by field (`Configuration` is not copyable) and compiles the library's sources plus `tests/<name>/main.cpp` into one executable, so a test reaches non-exported internals the way the library does. A new test is a directory, a `cfg.tests.push_back(LibraryTest("Name"));` and a comment above it saying what it covers and what it needs (a GPU, ray tracing, a compositor, a tool). A test that needs shaders adds them to `test.config.shaders` with the same `ShaderType` lines the examples use; a font or asset goes in `test.config.files`; a tool it runs goes in `requires_` (`{"tool:glslang", "tool:spirv-val"}`), which turns the test into a skip when the tool is missing. The declarations sit in the native branch of `project.cpp`: wasm runs no tests. Wayland-only tests sit in `if (!windows)`; `PointerWin32` is the one Win32 test, with `requires_ = {"tool:wine"}` when built on Linux.
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/<name>.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`, `WorkaroundNvidiaAS::Patch`, the font and input-field helpers). `PushConstantRewrite` and `UIFusedShader` compile GLSL with glslang and check the SPIR-V with spirv-val at runtime; `MouseScroll` and `PointerEnterLeave` drive the Wayland listener callbacks directly.
- 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`, 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) and check pixels.
Conventions the existing tests follow:
- Open `main.cpp` with a comment that states the behaviour under test, the failure it guards against and what it needs. `tests/HeapSlotReuse/main.cpp` and `tests/UIFallbackParity/main.cpp` show the shape.
- `Check(ok, what)` prints `PASS` or `FAIL` with the description and counts; `main` prints a summary and returns `EXIT_FAILURE` when any failed. Do not stop at the first failure unless later steps are meaningless without it.
- A validation check reads `Device::validationErrorCount` (the debug messenger only counts, it never aborts) and is gated on `std::getenv("CRAFTER_GRAPHICS_VALIDATION")`, since without the layer the count is always zero.
- Process-wide state is tested in a child process: `UIFallbackParity` starts itself twice with `--render <png>`, once with `CRAFTER_GRAPHICS_NO_DESCRIPTOR_HEAP` set, and compares the frames; `FeatureRequests` runs its Vulkan 1.3 case under `CRAFTER_GRAPHICS_NO_VULKAN_1_4` the same way.
- 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.
Cross targets: `crafter-build test --target=x86_64-w64-mingw32` needs `wine` on the sandbox's `PATH`, which the Flatpak does not have, and a MinGW import library for `vulkan-1`, which the toolchain does not ship; the Win32 test has been run by building an import library from the `vk*` imports with `llvm-dlltool` and running the exe 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.
## Code style
Nothing enforces it here, so match the surrounding code:
- The first two lines of every first-party source are `//SPDX-License-Identifier: LGPL-3.0-only` and `//SPDX-FileCopyrightText: Copyright (C) 2026 Catcrafts®`, no space after the slashes, then a blank line. Examples are `MIT`. Shaders carry the same two lines as `//` comments. Vendored files keep their own headers.
- `import std;` everywhere. Standard headers are never `#include`d; the Vulkan header, the Wayland and Win32 headers, `stb_truetype.h` and the Linux input headers go in the global module fragment after `module;`, each behind its backend's `#ifdef`. Tests may `#include <cstdlib>`, `<cstring>` or `<cstddef>` for the macros `import std;` does not provide.
- Four-space indentation in new code. A few files have tab-indented blocks (`Semaphores` in `Window.cppm`, parts of `Window.cpp` and `VulkanBuffer.cppm`); match the block you are in and do not reindent. Eight files have no final newline; leave that alone too. There is no `.editorconfig` or `.gitattributes`.
- `PascalCase` for types, functions and statics, `camelCase` for members and variables, a trailing `_` on private members in the newer classes (`heap_`, `shapedRuns_`), `k`-prefixed constants (`kUIFlagClip`, `kSwapchainWriterStages`, `kRTGeometryInstanceForceOpaque`). Wayland and Win32 callbacks keep the protocol's `snake_case` names. Vulkan structs are filled with designated initialisers.
- Fixed-width types (`std::uint32_t`; `std::uint16_t` for slot indices); `int` only at C boundaries. `std::println` and `std::format`; `std::runtime_error` with a message naming the file or the field for anything a user can get wrong, `std::invalid_argument` for a precondition.
- Comments explain why: the failure mode, the constraint, the rejected alternative, the issue number (`#118`) on the Forgejo tracker. The codebase carries long comments on purpose, and the recent commit bodies have the voice. Keep them, and write them for new non-obvious code.
- GLSL: `#version 460`, `#extension GL_GOOGLE_include_directive : enable`, `#include "ui-shared.glsl"` (crafter-build puts `shaders/` on the include path through `cfg.buildFiles`), the `UI_HEAP_*` macros instead of `descriptor_heap`, `std430` structs mirrored by a C++ struct with a `static_assert` on its size. Examples ship a `.wgsl` beside every `.glsl`.
## Design rules
These were decided deliberately. Do not undo them without the maintainer's say-so.
1. **The target picks the backend, and there is no abstraction layer.** `project.cpp` defines exactly one of `CRAFTER_GRAPHICS_WINDOW_WAYLAND`, `CRAFTER_GRAPHICS_WINDOW_WIN32` and `CRAFTER_GRAPHICS_WINDOW_DOM` from the target triple; there is no runtime backend switch and no X11. `:GraphicsTypes` holds `using` aliases, not virtual interfaces. Where the two backends' APIs differ, portable code `#ifdef`s at the call site, as the examples do, rather than the library inventing a shared signature that is a no-op on one side.
2. **Every partition exists on every target.** crafter-build's dependency scanner does not respect `#ifdef` around an `import :X;`, so a partition file is never left out of the list; its body is gated instead, and the umbrella imports everything unconditionally. Backend-specific implementation units are chosen in `project.cpp`.
3. **One binding model in two builds.** On the heap path pipelines are created with `VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT`, no pipeline layout, and push data through `vkCmdPushDataEXT`. On a driver without the heap the same slots live in one bindless descriptor set per frame (bindings 0 to 3: storage images, sampled images, samplers, storage buffers) and push constants go through `vkCmdPushConstants`. Slot handles, item structs, `UIDispatchHeader` and the `UIRenderer` API mean the same in both. Every compute kernel is compiled twice from one source (`CRAFTER_GRAPHICS_DESCRIPTOR_SETS` selects the `UI_HEAP_*` macros; the variant is `sets`), `ComputeShader::Load` picks the file, and `UIFallbackParity` proves the two render identically. Push blocks stay within 256 bytes. Ray tracing is heap-only.
4. **The UI contracts are frozen.** `UIDispatchHeader` is 48 bytes and starts every UI push block; `QuadItem` is 64, `CircleItem` 48, `ImageItem` 64 and `GlyphItem` 48 bytes, byte for byte with `ui-shared.glsl` under `std430`. They grow only through the reserved `flags` bits and `_pad`. `UIFusedHeader` (128 bytes) is a separate contract for `DispatchFused`.
5. **Three tiers, and the source is the customisation API.** Tier 1 is `ComputeShader` and `UIRenderer::Dispatch`; Tier 2 the standard kernels, item structs and `UIRenderer` helpers; Tier 3 stateless functions that append items to user-owned buffers and never dispatch. No library `Theme` type (colours come as `*Colors` aggregates), no widget tree, no layout engine beyond `Rect::SubRect` and `Inset`, no hit-testing, no focus management; state lives in user-owned PODs (`InputField` is the model). A component that does not fit is copied and edited, not hooked: no virtual overrides, no extension points. Additive changes (another kernel, another stateless component) are fine.
6. **Standard kernels dispatch one workgroup per 8x8 screen tile**, each thread owning one pixel and iterating every item in array order into a local accumulator, stored once. Array order is draw order and there is no inter-workgroup race. Never go back to a workgroup per item: Vulkan orders nothing between the workgroups of one dispatch, and overlapping items then race on the image (striped fills, stale pixels). The cost is `O(W*H*N)` and documented as such; the shared-memory tile cull in `ui-shared.glsl` must stay a strict superset of the per-pixel test.
7. **Three frames in flight, a fence per image, no steady-state wait-idle.** The owners of buffers, acceleration structures, shader modules and pipelines (`VulkanBuffer`, `Mesh`, `TlasWithBuffer`, `VulkanShader`, `ComputeShader`, `PipelineRTVulkan`) release them through `Device::EnqueueDeletion`, which has overloads for exactly those handle kinds; `ReclaimDeletions` frees an entry `framesInFlight` frames later and `DrainDeletions` after a wait-idle. Owners of handles the queue has no overload for wait for the GPU first and then destroy directly: `~DescriptorHeapVulkan` in descriptor-set mode, `~Window`, `~FontAtlas` through `ImageVulkan::Destroy`. A destructor never destroys a handle without one or the other: a frame in flight may still use it, and a static owner of a queued kind is destroyed after the driver has gone, which only the queue survives. `Clear()` is for a GPU known to be idle. The deletion queue, its mutex and `RenderingElement3D::tlases` are leaked on purpose for the same reason. Staged re-uploads go through the per-frame staging ring in `VulkanBuffer::Stage`.
8. **The library enables only the device features its own shaders use.** Everything an application's shaders need (`shaderInt16` for a `uint16_t` spec constant, fp16, 16-bit storage, `scalarBlockLayout`) is the application's to ask for through `Device::Initialize(required, optional)`, and `VulkanShader` refuses a module whose capabilities nothing enabled, naming the field. Do not add a feature to the `library` set for an example's convenience; have the example ask.
9. **Barriers are scoped and derived.** A `RenderPass` reports its `SwapchainStage()`; the frame loop's acquire, inter-pass and present barriers take their stage and access masks from the passes (`SwapchainStageUnion`, `SwapchainWriterAccess`, `BuildSwapchainInterPassBarrier`), masked by `Device::stageMask` so the RT stage never appears on a device without it. In the frame loop there is no `ALL_COMMANDS` and no queue-wide memory barrier, and a ranged flush rounds to `nonCoherentAtomSize`. The acceleration-structure builds and the GPU decompression do use global `VkMemoryBarrier`s on purpose (`BuildTlasWriteBarrier`, `BuildTlasReadBarrier`, `Decompress::DecompressOnGPU`): execution order alone never makes a build's writes visible to the shaders that read them. The window never clears the swapchain image (`TRANSFER_DST` is not on every surface); passes draw their own background.
10. **Driver workarounds are fenced and written to be deleted.** The NVIDIA descriptor-heap AS-read rewrite is one block in `ShaderVulkan.cppm` gated on `Device::workaroundDescriptorHeapAS`, with its push-data hook in `RTPass` and `ComputeShader`, a kill switch (`CRAFTER_GRAPHICS_NO_AS_WORKAROUND`), a probe test and a status note in the README and `examples/VulkanTriangle/README.md`. A new workaround gets the same: a `Device` flag, one fenced block, a probe, a note. A fault the library cannot hide (dynamic descriptor indexing in hit shaders) is documented with the consumer-side pattern instead (`examples/Sponza/README.md`).
11. **Prefer the clean breaking change.** When a fix can be a backward-compatible patch or a cleaner API change, make the API change and update the examples and tests in the same commit; the consumers are in this tree or the maintainer's. Do not stage a compat shim "for now".
12. **`lib/` is vendored and untouched.** libevdev is built as its own static library configuration so it compiles once, not once per test; its include path and the kernel headers are absolute because this `project.cpp` is also loaded from a consumer's directory.
## Commit and PR guidelines
- Branch from `master`; PRs go to `master` on `forgejo.catcrafts.net/Catcrafts/Crafter.Graphics`. Agent branches are `claude/<topic>`, issue work `claude/issue-<n>`. The fetch URL is HTTPS and the push URL SSH (`forgejo@forgejo.catcrafts.net:Catcrafts/Crafter.Graphics.git`); the server refuses HTTPS pushes.
- Subject: `type(scope): imperative summary`, with `feat`, `fix`, `perf`, `refactor`, `docs` or `test` and the area as the scope (`device`, `rt`, `ui`, `window`, `wayland`, `win32`, `webgpu`, `font-atlas`, `compute`, `gamepad`). The body is grouped under the components it touches (`Device`, `Shaders`, `Window`, `Tests`) and says what was wrong, why, and why the fix has this shape; `git log -3` shows the voice. Add a `Co-Authored-By:` trailer when an agent wrote it.
- Update the README (and the example README that holds the investigation) in the same change when user-visible behaviour changes, and `AGENTS.consumer.md` when a convention, a signature, a contract or an environment variable consumers lean on changes. Consumers pick the new notes up on their next build.
- Crafter.Event, Crafter.Math and Crafter.Asset are pulled unpinned, so a change that needs one of them lands there first and is built here with `--local` meanwhile. Pushing is per repository: land the sibling, then this one.
- Never commit `bin/`, `build/` or `pipeline_cache.bin`, and never edit `lib/`.
- A second worktree of this repository may sit beside it (`git worktree list`); its branch shows with a `+` in `git branch`. Leave it to whoever opened it.
## Security considerations
- `project.cpp` is compiled and run by crafter-build; it clones and pulls three sibling repositories from their default branch and two Khronos repositories from GitHub on every invocation, and the Sponza example `GitFetch`es a 280 MB asset repository (pinned to a commit). Pin a `commit` in a `GitProject` spec when that matters.
- The validation layer, GPU-assisted validation included, is a developer tool and off by default; `CRAFTER_GRAPHICS_VALIDATION` turns it on. GPU-AV has crashed a vendor SPIR-V compiler on the heap pipelines before, which is why it is opt-in.
- Shaders are read from disk relative to the working directory and checked for the SPIR-V magic and word alignment before they reach the driver; the NVIDIA rewrite edits SPIR-V in memory and `PushConstantRewrite` validates its output with spirv-val.
- `pipeline_cache.bin` is written to the working directory at exit and read back only when its header matches the device; it is gitignored here and should be in a consumer.
- The environment switches (`CRAFTER_GRAPHICS_VALIDATION`, `CRAFTER_GRAPHICS_DEVICE`, `CRAFTER_GRAPHICS_NO_DESCRIPTOR_HEAP`, `CRAFTER_GRAPHICS_NO_AS_WORKAROUND`, `CRAFTER_GRAPHICS_NO_VULKAN_1_4`, `CRAFTER_PRESENT_IMMEDIATE`) are read with `std::getenv` in `Device.cpp` and `Window.cpp`. A new one takes the `CRAFTER_GRAPHICS_` prefix, means "set to any value", and is listed in the README and in `AGENTS.consumer.md`.
- `dom-env.js` and `dom-webgpu.js` are shipped to browsers as they are. `dom-webgpu.js` replaces `document.body` with an error when `navigator.gpu` is absent, which is why `--no-webgpu` exists for DOM-only pages.
- Tests write only under `fs::temp_directory_path()` and their own bin directory.
## Gotchas
- A new partition is three edits: the interface file, the `ifaces` array in `project.cpp` (a `std::array` with a literal size, so bump it) and the umbrella `Crafter.Graphics.cppm`. An implementation unit goes in `impls` (native), `domImpls` (DOM) or both, again with literal sizes; a test's sources are built from `impls`, so a native unit missing there fails every test link.
- `Device::Initialize` on the Wayland backend connects to the display and throws without one; "headless" tests still need a compositor.
- Buffer slot indices: `AllocateBufferSlots` returns an index relative to the heap's buffer region, while GLSL `descriptor_heap` SSBO arrays index in buffer-descriptor units from heap byte 0. `BufferSlot` converts to `bufferStartElement + raw`; raw code (the RT examples) passes `heap.bufferStartElement` as a `uint16_t` spec constant, which needs `shaderInt16`. In descriptor-set mode `bufferStartElement` is 0. Image and sampler slots are used as they are. A wrong index reads zeros, which looks like "every item field is zero".
- `UIRenderer::RegisterBuffer` refuses a buffer created without `VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT`. Freed slots are handed out again at once, so a slot may only be freed when no frame in flight reads it (after `vkQueueWaitIdle`, or at a scene switch). In descriptor-set mode a slot write destroys the previous view immediately.
- The swapchain image stays in `VK_IMAGE_LAYOUT_GENERAL` across the pass list and is never cleared; `Window::clearColor` is unused. `swapchainUsage` says whether `TRANSFER_DST` was available. `onResize` fires after a swapchain rebuild; anything holding descriptors of `window.imageViews[]` rewrites them there (`UIRenderer` does its own).
- `Device::frameCounter` is one clock for every window, so the deletion queue and the staging rings assume one rendering window.
- `ImageVulkan::Create`'s default `consumerStage` is the RT stage, corrected to compute on a device without ray tracing; the font atlas passes compute explicitly. `VulkanTransition`'s `GENERAL` mapping does the same.
- `UIRenderer::ShapeText` and `FontAtlas` key their caches on the `Font*`: call `UIRenderer::InvalidateFont` before destroying a `Font`, or a new one at the same address aliases stale runs.
- `RenderingElement3D::elements` stores raw pointers; a `std::vector<RenderingElement3D>` that grows after `Add` dangles them (Sponza reserves up front). `BuildTLAS` on an index whose earlier submission has not completed overwrites a TLAS in use; building `tlases[window.currentBuffer]` from `onUpdate` is safe, and `tlases[i].address` changes when a TLAS outgrows its storage.
- `PipelineRTVulkan::Init` defaults `maxRecursionDepth` to 1; a closest-hit that traces needs 2, and tracing deeper than the pipeline allows loses the device.
- `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`.
- 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.
<!-- crafter-build: dependency notes begin -->
<!-- Written by crafter-build from the AGENTS.consumer.md each dependency ships, and rewritten on every build. Edit those files, not this section; cfg.dependencyNotes in project.cpp picks which dependencies are included. -->
<!-- from Crafter.Math -->
## Crafter.Math
The Crafter engine's math library: a C++ module with no implementation units and no dependencies, so everything is a template or an inline body compiled into your translation units with your build's `--target` and `--march`. Source and issues: https://forgejo.catcrafts.net/Catcrafts/Crafter.Math. Namespace `Crafter`. The README there documents the whole surface; these notes are what an agent working in a consumer needs.
### Which type
**`Vector<T, Len, Alignment = 0>`** is plain data, for vertex and uniform structs, events, anything stored or passed around. `.x .y .z .w` (and `.r .g .b .a`) alias `v[]`; constructors exist for `Len` 2, 3 and 4 (`Vector<float, 2>{x, y}`); `+ - * /` with vectors and scalars, `+=` and friends, `==`, unary `-`; statics `Dot(a, b)`, `Cross(a, b)` (`Len == 3`), `Normalize(a)`; members `Length()`, `LengthSq()`, `Normalize()`, `Invert()`. `std::format("{}", v)` prints `{x,y,z}`.
- `Alignment` is the number of stored elements. `0` means `Len`, naturally aligned: `Vector<float, 3>` is 12 bytes aligned 4. Any `N >= Len` stores `T v[N]` and marks the struct `__attribute__((packed))`, alignment 1: `Vector<float, 3, 3>` is 12 bytes aligned 1 (the vertex type Crafter.Graphics uses), `Vector<float, 3, 4>` is 16 bytes, a `vec4` slot. Conversion between alignments is implicit.
- Arithmetic between two lengths covers the shorter one and leaves the other lanes of the result uninitialised, and `==` compares the shorter. The member `Normalize()` divides without a zero check; the static `Normalize(a)` checks.
**`VectorF32<Len, Packing>`** is one SIMD register holding `Packing` vectors of `Len` floats, for math-heavy loops. Construct from a `const float*` or broadcast one `float`; `Load(ptr)`, `Store(ptr)`, and `Store<float>()` returning a `std::array`. Operators as above; statics `Cross` (`Len == 3`), `MulitplyAdd(a, b, c)` (`a*b + c`), `MulitplySub(a, b, c)` (`a*b - c`), `Rotate(v, q)` and `RotatePivot(v, q, pivot)` by a `VectorF32<4, Packing>` quaternion, `QuanternionFromEuler(halfAngles)`, `Blend<{{bools}}>(a, b)`; members `Sin()`, `Cos()`, `SinCos()`, `Shuffle<{{i, j, k}}>()`, `Negate<{{bools}}>()`, `ExtractLo<N>()`. Inside a template write `v.template Store<float>()`.
- `Load`, `Store` and the pointer constructor move the whole register, `AlignmentElement` floats (4, 8 or 16), not `Len * Packing`. Give them a buffer that big (`alignas(16) float buf[4] = { x, y, z, 0 };` is the idiom in the tests), zero the padding lanes on the way in and ignore them on the way out. The pointer need not be aligned.
- `Dot`, `Length`, `LengthSq` and `Normalize` are batched with the arity fixed at compile time: `Normalize`, `Length` and `LengthSq` take exactly `BatchSize` vectors, `Dot` takes `2 * BatchSize` (a0, b0, a1, b1, ...). They return `VectorF32<1, Packing * BatchSize>` with one lane per input vector (`Normalize` returns `std::array<VectorF32<Len, Packing>, BatchSize>`). There is no two-argument `Dot`: for one product pass the same pair `BatchSize` times and read lane 0, as `MatrixRowMajor::operator*` does, or use `Vector<float, 3>::Dot`.
- `BatchSize` is `AlignmentElement / Packing`, 4 for every `VectorF32<Len, 1>`. `OptimalPacking` is how many `Len`-vectors fill one register (table below). Shapes convert implicitly, copying the lanes they share.
- The x86 backend lacks `const` on the scalar operators, `Sin`, `Cos`, `SinCos` and `Shuffle`: call them on non-const objects.
**`VectorF16<Len, Packing>`** is the same surface for `_Float16` where the target has native half arithmetic: x86 with AVX512-FP16, AArch64 with ARMv8.2-FP16. Everywhere else, including every x86 tier below `sapphirerapids` and all of RISC-V, it is `using VectorF16 = VectorF32`, converting on `Load` and `Store` from `_Float16*`. Never overload a function on `VectorF16` and `VectorF32` both, and never assume its lanes are halves.
**`MatrixRowMajor<float, 4, RowSize, 1>`** is four-column and row-major, each row a `VectorF32<4, 1>` in the public `rows[]`; constructors take the 12 or 16 floats row by row. Only `T = float`, `CollumSize = 4`, `Repeats = 1` exist.
- `<float, 4, 3, 1>` (48 bytes) is the column-vector shape: `M * VectorF32<3, 1>` applies it with an implicit `w = 1`, `TransformNormal(v)` with `w = 0`; translation sits in column 3. `Store(float*)` writes the 12 floats row-major, which is `VkTransformMatrixKHR`: `m.Store(reinterpret_cast<float*>(instance.transform.matrix))`.
- `<float, 4, 4, 1>` (64 bytes) is the row-vector shape for the GPU side (`v * M`, translation in row 3). There is no matrix-times-vector operator for it in C++. `operator*` between two 4x4s, `Transpose()`, `LookTo(eye, direction, up)` and `LookAt(eye, focus, up)` are 4x4 only.
- Both shapes have `Identity()`, `Scaling(x, y, z)`, `Translation(x, y, z)` and `Rotation(pitch, yaw, roll)` in radians (each also from a `VectorF32<3, 1>`), `Rotation(Vector<float, 4> quaternion)`, `TRS(t, q, s)` (scale, then rotate, then translate: a glTF node), `InverseAffine()` (any rotation, scale and shear plus a translation, never a projection) and `Decompose()` returning `(t, q, s)`, with a mirroring expressed as a negative X scale. Each means the same transform in both shapes. `std::format` prints both.
**Quaternions** are `Vector<T, 4>` in `(x, y, z, w)` order, unit, `w` the scalar part, applied by `Vector<T, 3>::Rotate(v, q)` and `RotatePivot(v, q, pivot)`; identity is `(0, 0, 0, 1)`. The functions are statics of `Vector<T, 4>` and are spelled `Quanternion`: `QuanternionMultiply(a, b)` is the Hamilton product, b first then a (`Rotate(v, a*b) == Rotate(Rotate(v, b), a)`); `QuanternionConjugate(q)` is the inverse rotation; `QuanternionFromAxisAngle(unitAxis, radians)` is right-handed (+90° about Z takes X to Y); `QuanternionFromUnitVectors(from, to)` is the shortest arc, defined for opposite directions too; `QuanternionFromBasis(right, up, forward)`; `QuanternionFromEuler(rollHalf, pitchHalf, yawHalf)` takes **half** angles; `QuanternionNlerp(a, b, t)` and `QuanternionSlerp(a, b, t)` take the shorter arc and return unit quaternions.
**Intersection tests** (`IntersectionTestRayTriangle`, `IntersectionTestRaySphere`, `IntersectionTestRayOrientedBox`, `IntersectionTestSphereOrientedBox`, `IntersectionTestOrientedBoxOrientedBox`, with `PackedOBBs<Packing>` and `GetOBBCorners`) are batched over `Packing * BatchSize` primitives per call; `Packing` defaults to `VectorF32<3, 1>::OptimalPacking` (the ray-box test to the smaller of the 3- and 4-vector packings). A vertex argument is `std::array<VectorF32<3, Packing>, BatchSize>`: element `b`, lanes `s*3 .. s*3+2`, holds primitive `b * Packing + s`. The result is a `VectorF32<1, Total>` with lane `i` for primitive `i`: the ray parameter `t`, or `std::numeric_limits<float>::max()` for a miss (the two box-overlap tests return `0` for a hit). The triangle test culls back faces; the sphere test returns the near root, so a ray starting inside misses. `tests/Intersection/main.cpp` in the Crafter.Math repository has `PackVec3`, `PackVec4` and `PackScalars` helpers worth copying.
### Targets and SIMD
| Backend | Register | `MaxElement` (floats) | `VectorF32<3,1>::OptimalPacking` | `BatchSize` of `VectorF32<3, Opt>` |
|---|---|---|---|---|
| x86 AVX-512 (`x86-64-v4`, `sapphirerapids`) | 64 B | 16 | 5 | 3 (15 lanes) |
| x86 AVX2 (`x86-64-v3`), and the scalar fallback | 32 B | 8 | 2 | 4 (8 lanes) |
| AArch64 NEON, WebAssembly SIMD128 | 16 B | 4 | 1 | 4 |
| RISC-V RVV | 16, 32 or 64 B from the `Zvl*` in the march | 4, 8 or 16 | | |
- The backend and register width are fixed when the module is compiled, from predefined macros, so from the `--march` your build uses. crafter-build's default is `native`: fine for a developer build, wrong for a binary that ships. `--march=x86-64-v3` is the lowest x86 tier the library is tested on (its SSE paths use FMA); the `sapphirerapids` tier runs only on CPUs with AVX512-FP16.
- Tested as of 2026-10-31: all three x86 tiers on Linux and on Windows (`x86_64-w64-mingw32` under wine); AArch64 `armv8-a` and `armv8.2-a+fp16` under qemu, except Crafter.Math's own Matrix test, which does not compile for AArch64 with clang 23 (`arm_neon.h` wrapper functions defined twice across module units; it bites a unit that emits `VectorF32<3, 1>` arithmetic before a `VectorF32<3, 1>` `Sin`, `Cos` or `SinCos`, as that test does through `MatrixRowMajor::Rotation(pitch, yaw, roll)`). If an AArch64 build of yours fails with `definition with same mangled name` in `arm_neon.h`, it is that issue. RISC-V has never been built. WebAssembly builds (crafter-build adds `-msimd128`; flags you add to your own units must match) and has no tests.
### Spellings
- `Quanternion*`, `MulitplyAdd`, `MulitplySub`, `CollumSize`, `Aligment` are the real identifiers. Don't correct them in a consumer; a rename is an upstream change that updates every consumer at once.
<!-- crafter-build: dependency notes end -->