diff --git a/AGENTS.consumer.md b/AGENTS.consumer.md new file mode 100644 index 0000000..3cab840 --- /dev/null +++ b/AGENTS.consumer.md @@ -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` from Crafter.Event (`import Crafter.Event;`): `EventListener 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 (`.spv` from `.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=` 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`: `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` 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::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` (`Image2D` 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` 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` 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(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 `.spv` and `.sets.spv`. `ComputeShader::Load(".spv", heap)` picks the build the device needs and throws naming a missing file. `examples/CustomShader/project.cpp` has the lines. +- `UIDispatchHeader` (48 bytes; in C++ `outImage, itemBuffer, surfaceWidth, surfaceHeight, clipX, clipY, clipW, clipH, itemCount, frameIdx, flags, _pad`, and in GLSL the same bytes as `uint outImage, itemBuffer; uvec2 surfaceSize; vec4 clipRectPx; uint itemCount, frameIdx, flags, _pad`) must be the first member of the push block; `ui.FillHeader(slot, count, clipRectPx, flags)` fills it. The high bit of `flags` (`kUIFlagClip`) is the renderer's, the rest are yours. Keep the push block a multiple of 4 bytes and within 256; the descriptor-set path throws past the device's `maxPushConstantsSize`. +- `ui.Dispatch(cmd, shader, &pc, sizeof(pc), ceil(W / 8), ceil(H / 8))` binds, pushes and dispatches, with a `SHADER_WRITE -> SHADER_READ | SHADER_WRITE` barrier on the swapchain image before every dispatch but the frame's first, so a custom shader sees what the dispatch before it drew. `ComputeShader::Dispatch(cmd, push, bytes, gx, gy, gz, tlasAddress = 0)` is the same without the barrier; a compute shader that ray-queries the TLAS passes `RenderingElement3D::tlases[frameIdx].address`, for the NVIDIA workaround below. +- 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> vertices, std::span 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::Store(reinterpret_cast(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` 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`: whole detents, reinterpret the payload as `int32`, positive is wheel down, every backend normalised to one per notch), `onRawKeyDown`, `onRawKeyHold`, `onRawKeyUp` (`Event`, raw platform codes; `heldKeys`), `onTextInput` (`Event`, UTF-8; repeats are synthesised on Wayland), `onPreedit` and `onTextInputDelete` (input methods), `onTouchDown`, `onTouchMove`, `onTouchUp` (`Event`; the first finger also drives the pointer events), `onClose`, `onResize`, `onBeforeUpdate`, `onUpdate` (`Event`: `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`). + +### 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::Create(count)`, `FlushDevice()` without a command buffer (`queue.writeBuffer`), `FlushDeviceRange`, `FlushDeviceSlot`, and `EnqueueReadback` with `PollReadback` for GPU to CPU. `Image2D::Create(w, h)` is `rgba8unorm`, `Update(const CompressedTextureAsset&)` decompresses on the CPU, `AllocateSlot(heap)` registers it; `Image2DArray` 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=` | 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. diff --git a/AGENTS.md b/AGENTS.md new file mode 100644 index 0000000..4db9e25 --- /dev/null +++ b/AGENTS.md @@ -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-.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-.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//` | 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//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----/` 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:` 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/ && 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//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/.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 `, 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 `, `` or `` 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/`, issue work `claude/issue-`. 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` 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.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`** 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{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` is 12 bytes aligned 4. Any `N >= Len` stores `T v[N]` and marks the struct `__attribute__((packed))`, alignment 1: `Vector` is 12 bytes aligned 1 (the vertex type Crafter.Graphics uses), `Vector` 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`** 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()` 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()`. Inside a template write `v.template Store()`. + +- `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, 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::Dot`. +- `BatchSize` is `AlignmentElement / Packing`, 4 for every `VectorF32`. `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`** 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`** 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. + +- `` (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(instance.transform.matrix))`. +- `` (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 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` in `(x, y, z, w)` order, unit, `w` the scalar part, applied by `Vector::Rotate(v, q)` and `RotatePivot(v, q, pivot)`; identity is `(0, 0, 0, 1)`. The functions are statics of `Vector` 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` 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, 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::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. + +