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 |
| `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`, `DescriptorHeapVulkan::ComputeLayout`, the font and input-field helpers). `PushConstantRewrite` and `UIFusedShader` compile GLSL with glslang and check the SPIR-V with spirv-val at runtime, and `ExampleShaders` compiles every GLSL shader under `examples/` the way crafter-build would, the descriptor-set build included, so a shader contract change that breaks an example fails here without building the example; `MouseScroll` and `PointerEnterLeave` drive the Wayland listener callbacks directly.
- A headless device: `Device::Initialize()` and no `Window` (`HeapSlotReuse`, `DeferredDeletion`, `MeshLifetime`, `ImageStagingRelease`, `BLASBuildOptions`, the `TLAS*` tests, ...). Headless still means a Wayland display on Linux. Most of the RT ones check `Device::rayTracingSupported` and return 77 without it; `BLASBuildOptions`, `TLASInstanceDirtyTracking` and `MeshDecompressStagingRelease` do not, and crash on a GPU without ray tracing (`Device::RequireRayTracing` throws out of `main`). `DeviceSelection` skips with one GPU.
- A window and a live compositor: `FrameLoopSync`, `SetupCmdBufferReuse`, `UIFallbackParity`, `SwapchainTransferDst`, `WindowUpdateClock`, `HeapASRead`, `HeapDynamicIndexHit`, `RTPassPushData`, `WindowOptionalGlobals` (hides xdg-decoration, viewporter and fractional scale from the window in turn), 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.
- 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.
| 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.
