Crafter.Graphics/AGENTS.md

AGENTS.md

Guidance for coding agents and contributors working on Crafter.Graphics itself. The README is the user-facing manual, the example READMEs hold the driver investigations, and 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 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 imports 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, 3.0.0 or later (project.cpp uses its ShellQuote), 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 pulls them on every invocation, and clones the two Khronos header repositories, and for a Win32 target Vulkan-Loader as well, 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; for a Win32 target it links the loader through an import library that project.cpp makes with llvm-dlltool from Vulkan-Loader's loader/vulkan-1.def, into build/vulkan-1-<target>/, since the MinGW toolchains ship none. 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), ray tracing and swapchain_maintenance1 (with the instance's surface_maintenance1, for per-present-mode image counts) 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 (an import library made from Vulkan-Loader's loader/vulkan-1.def), kernel32, user32, gdi32, runtimeobject, uuid; 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

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 --target=x86_64-w64-mingw32 --runner=cmd:true Win32Link   # links a Win32 executable, without Wine
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 test --target=x86_64-w64-mingw32 --runner=cmd:true Win32Link still links when you touched Window, Device, Gamepad, Clipboard or the link flags in project.cpp.

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); the Win32 tests, PointerWin32 and TextWin32, sit in if (windows), with requires_ = {"tool:wine"} when built on Linux. Win32Link beside them has no requires_, since its link is the test: where there is no Wine it runs as --runner=cmd:true, which builds it and runs nothing. A requires_ would skip the build as well, and a test cannot pick its own runner, as the harness derives every test's from its target.

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, Window::SwapchainMinImageCount, WorkaroundNvidiaAS::Patch, DescriptorHeapVulkan::ComputeLayout, FontAtlas::Packer, the font and input-field helpers). PushConstantRewrite and UIFusedShader compile GLSL with glslang and check the SPIR-V with spirv-val at runtime, and ExampleShaders compiles every GLSL shader under examples/ the way crafter-build would, the descriptor-set build included, so a shader contract change that breaks an example fails here without building the example; FontAtlasDecode reads ui-shared.glsl and dom-webgpu.js and holds their copies of the SDF encoding constants to FontAtlas's; MouseScroll and PointerEnterLeave drive the Wayland listener callbacks directly, and TitleWayland reads the requests a window sends off a wl_display connected to a socketpair.
• A headless device: Device::Initialize() and no Window (HeapSlotReuse, DeferredDeletion, MeshLifetime, ImageStagingRelease, BLASBuildOptions, the TLAS* tests, ...). Headless still means a Wayland display on Linux. Most of the RT ones check Device::rayTracingSupported and return 77 without it; BLASBuildOptions, TLASInstanceDirtyTracking and MeshDecompressStagingRelease do not, and crash on a GPU without ray tracing (Device::RequireRayTracing throws out of main). DeviceSelection skips with one GPU.
• A window and a live compositor: FrameLoopSync, SetupCmdBufferReuse, UIFallbackParity, SwapchainTransferDst, WindowUpdateClock, HeapASRead, HeapDynamicIndexHit, RTPassPushData, WindowOptionalGlobals (hides xdg-decoration, viewporter and fractional scale from the window in turn), SwapchainImageCount (every present mode, and swapchains with more images than frames in flight, in child processes on llvmpipe where the host's GPU never makes one), TextAntialiasing (text through DispatchText and DispatchFused against stb_truetype's own antialiased rasterizer, at 14 and 64 px and stretched, plus a glyph mirrored by its UVs and empty texels shrunk below 4 px), and FeatureRequests for its retry case. They render, read the frame back with Window::CaptureNextFrame (a PNG, decoded with a private STB_IMAGE_STATIC copy of stb_image) or a transfer pass of their own, and check pixels.

Conventions the existing tests follow:

• 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. TextAntialiasing compares the two for text, and keeps the text's edge ramp at one screen pixel; for quads, circles and images no test does (UIFusedShader only pins the push-constant layout): check those by eye, or add the comparison.

Cross targets: the Win32 tests run under Wine from the native crafter-build on the host, which has wine on its PATH and finds the host's mingw-w64 itself: WINEPREFIX=<scratch prefix> crafter-build test --target=x86_64-w64-mingw32 PointerWin32 TextWin32 Win32Link. The Flatpak has no wine, so there PointerWin32 and TextWin32 skip and only Win32Link's link is checked. Under the Flatpak a Win32 executable also links only with tests/Win32Link/libc-single-threaded.c. The MinGW sysroot the Flatpak picks is the host's /usr, which puts glibc's headers on the MinGW include path, and libstdc++'s shared_ptr, inlined from nlohmann/json in Crafter.Asset's GLB loader, then reads glibc's __libc_single_threaded (Crafter.Build #43). 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 #included; 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 #ifdefs 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 frame, no steady-state wait-idle. Frames in flight are not swapchain images: the swapchain has as many images as its present mode needs, what lives per frame (command buffer, fence, acquire semaphore, heap copy, TLAS) is keyed by Window::currentFrame and what lives per image by currentImage, and the image a frame draws into reaches shaders through the heap's swapchain slots (DescriptorHeapVulkan::AddSwapchainSlot, #160). 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 VkMemoryBarriers 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, for a Win32 target also clones and pulls Vulkan-Loader, unpinned, and runs llvm-dlltool on its loader/vulkan-1.def when the import library under build/ was made from another .def or command, and the Sponza example GitFetches 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; slots listed with AddSwapchainSlot follow by themselves, and anything else holding the old images is made again there. A rebuild asked for during a frame (Resize or SetPresentMode from a listener) waits for that frame's present, since the rest of the frame indexes the swapchain it acquired from. RenderPass::Record's frameIdx is the frame slot, not the image.
• Device::frameCounter is one clock for every window, so the deletion queue and the staging rings assume one rendering window. The font atlas evicts on it too, on both backends: without a Window advancing it nothing is ever evicted, and a full atlas leaves new glyphs blank, so a headless test steps it itself (FontAtlasEviction).
• 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, which clears both, before destroying a Font, or a new one at the same address aliases stale runs and glyphs. The atlas also keeps the Font* of a glyph waiting for a cell until the next frame's first call, where it rasterises from it.
• 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.currentFrame] 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.
• The text kernels decode FontAtlas's SDF in one place, ui-shared.glsl (UI_SDF_ON_EDGE, UI_SDF_PIXEL_DIST_SCALE, uiGlyphRamp, uiGlyphCoverage), and the WGSL text kernel in dom-webgpu.js repeats it; FontAtlasDecode holds both copies of the constants to FontAtlas's, and nothing renders the WGSL, so check a change to it by hand (examples/HelloUI built with --target=wasm32-wasip1, in a browser, or the kernel run through Deno's WebGPU). The WGSL lives in String.raw literals: a backtick in a WGSL comment ends the literal and breaks the whole backend, which only node --check additional/dom-webgpu.js notices.
• Crafter.Math's misspelled identifiers (Quanternion*, MulitplyAdd, CollumSize) are its API; do not correct them here.
• Dependency notes: crafter-build 3.0.0 or newer keeps a section at the end of this file holding the AGENTS.consumer.md of every dependency that ships one (2.2.0 leaves the file alone), rewritten on every crafter-build and crafter-build test. As of that date Crafter.Math's notes are on its master, Crafter.Asset's only on its claude/glb-loader branch, and Crafter.Event has none, so a plain build adds Math's and a --local build whatever the sibling checkouts hold. Edits inside the section's markers are lost on the next build; a change to those notes belongs in their repositories. Commit the section when it changes.

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.